{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "<a href=\"https://www.nvidia.com/dli\"> <img src=\"images/DLI Header.png\" alt=\"标题\" style=\"width: 400px;\"/> </a>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# 基于 Numba 的 CUDA Python 编程简介\n",
    "\n",
    "**[CUDA](https://en.wikipedia.org/wiki/CUDA)** 计算平台可以让开发者在 NVIDIA GPU 上实现大规模并行计算，从而加速应用程序的执行，节省宝贵的时间资源。\n",
    "\n",
    "**[Numba](http://numba.pydata.org/)** 是一款即时(JIT) Python 函数编译器，提供简洁的 API 用于加速 Python 函数运行。对希望使用 GPU 加速其应用又不想写 C/C++ 代码的同学而言，Numba 是个极具吸引力的选择，而对日常使用 NumPy 执行大量计算任务的开发者来说，Numba 更是不二之选。Numba 可用于加速 CPU 上运行的 Python 函数，同样也适用于在 NVIDIA GPU 进行应用加速。**本课程侧重讲解使用 Numba 在 GPU 上加速 Python 应用程序所需的基础技能。**"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 课程结构\n",
    "\n",
    "本课程分为以下**三个**部分：\n",
    "\n",
    "- _基于 Numba 的 CUDA Python 编程简介_\n",
    "- _使用 Numba 在 Python 中编写自定义的 CUDA 核函数_\n",
    "- _使用 Numba 实现 CUDA Python 的多维网格和共享内存_\n",
    "\n",
    "每节均包含一项最终评估，成功完成该评估后，我们将为您授予该课程的能力证书。每节另包含高阶资料附录，感兴趣者可借此深入学习。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 基于 Numba 的 CUDA Python 编程简介\n",
    "\n",
    "在第一部分中，您将首先学习如何使用 Numba 编译 CPU 函数，并且了解 Numba 编译器内部的工作过程。然后，您将继续学习使用 GPU 加速逐元素运算的 NumPy 数组函数，以及在 CPU 和 GPU 间高效传输数据的技术。\n",
    "\n",
    "在第一个部分结束后，您将能使用 GPU 对 NumPy 数组上的执行元素级运算的Python代码进行加速。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 使用 Numba 在 Python 中编写自定义 CUDA 核函数\n",
    "\n",
    "在第二部分中，您将进一步拓展能力，通过编写自定义 CUDA 核函数，实现在 GPU 上并行执行任意的而不仅仅是基于元素的计算密集型函数。为达成学习目标，您还需了解 NVIDIA GPU 如何并行执行代码，并学习几个基础的并行编程技巧，包括如何协调线程工作以及处理竞争条件(racing condition)。您还会学习在 GPU 上调试代码的技巧。\n",
    "\n",
    "第二部分学习结束后，您将能使用 GPU 加速各类处理一维数据的计算密集型函数。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 基于 Numba 控制 CUDA 的多维网格和共享内存\n",
    "\n",
    "在第三部分中，您将开始并行处理二维数据，并将学习如何使用 GPU 上的称为共享内存的片上存储空间。\n",
    "\n",
    "第三部分学习结束后，您将能使用 Numba 对一维和二维数据编写经 GPU 加速的 Python 代码。同时，您还可使用几种最为重要的优化策略，以编写稳定而快速的 GPU 程序。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 课程预备知识\n",
    "\n",
    "* 能够编写 Python程序，特别是会编写和调用函数，处理变量、循环、控制逻辑，导入和使用工具包。\n",
    "* 熟悉 Python 的 NumPy库。如果您从未用过 NumPy 但熟悉 Python，那么您会发现在本课中使用的 NumPy 功能是简单且直观的。对于需要阐明的知识点，我们会提供注释和链接。\n",
    "* 大致了解某些计算机科学的术语，如内存分配、数据类型、延迟和处理内核。\n",
    "* 基本理解向量、矩阵和矩阵乘法的概念。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 本节目标\n",
    "\n",
    "完成本节内容的学习后，您将能够：\n",
    "\n",
    "- 使用 Numba 编译运行在 CPU 上的 Python 函数。\n",
    "- 了解 Numba 如何编译 Python 函数。\n",
    "- 使用 GPU 加速 NumPy 通用函数 (ufunc)。\n",
    "- 使用 GPU 加速自定义的数组处理函数。\n",
    "- 优化 CPU 主机和 GPU 设备之间的数据传输。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 什么是 Numba？\n",
    "\n",
    "Numba 是一个**类型专用**的**即时函数编译器**，用于为 CPU 或 GPU 加速**以数值计算为主的** Python 函数。此定义很长，下面就让我们逐一解析这些术语：\n",
    "\n",
    "* **函数编译器**：Numba 用于编译 Python 函数，而非整个应用程序，它也不是函数的一部分。Numba 不会取代 Python 解释器，而仅作为另一个 Python 模块，将普通函数转化为执行速度更快的函数（通常情况下）。\n",
    "* **类型专用**：Numba 可为您当前使用的特定数据类型生成专门的执行代码，从而加速函数运行。Python函数被设计为对通用数据类型进行操作，这为其带来了极大的灵活性，但也严重拖慢了运行速度。实际上，您只会调用具有少量参数类型的函数，因此Numba会为每种类型的集合生成快速实现。\n",
    "* **即时**：在函数首次被调用时，Numba 会对它们进行转换，这样可以确保编译器知道您将使用的参数类型。这也使得Numba可以像传统应用程序一样轻松地在Jupyter笔记本中交互使用。\n",
    "* **以数字计算为主**：Numba 目前以处理数值型数据类型为主，如 `int`、`float` 和 `complex`。字符串处理支持极为受限，且许多字符串处理函数还无法在 GPU 上获得有效加速。若要借助 Numba 获得最佳加速效果，您可能需要搭配使用 NumPy 数组。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Numba 使用要求\n",
    "\n",
    "Numba 支持各类操作系统：\n",
    "\n",
    "* Windows 7 及以上版本，32 和 64 位\n",
    "* macOS 10.9 及以上版本，64 位\n",
    "* Linux（RHEL 5 及以上的大多数版本），32 和 64 位\n",
    "\n",
    "Numba 还支持多种 Python 版本：\n",
    "\n",
    "* Python 2.7、3.4 至 3.6 版本\n",
    "* NumPy 1.10 及以上版本\n",
    "\n",
    "此外，其还支持各类硬件：\n",
    "\n",
    "* x86、x86_64 及 AMD64 CPU\n",
    "* NVIDIA CUDA GPU（计算能力 3.0 及更高级别，CUDA 8.0 及以上版本）\n",
    "* AMD GPU（实验性补丁）\n",
    "* ARM（实验性补丁）\n",
    "\n",
    "本课程将使用 Linux 64 位和 CUDA 9 版本。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 注解：CUDA C/C++、Numba 与 PyCUDA 对比\n",
    "\n",
    "Numba 绝非 CUDA 编程的唯一手段。目前，最常见的 CUDA 编程方法便是使用 CUDA C/C++ 语言扩展程序。对于 Python 而言，[PyCUDA](https://documen.tician.de/pycuda/) 是除 Numba 以外可实现 GPU 加速 Python 代码的另一种选择。本课程将着重介绍 Numba。不过作为背景知识，我们也有必要在开讲前简单对比一下刚刚提到的三个选项。\n",
    "\n",
    "**CUDA C/C++**：\n",
    "- 最常用、高效且最灵活的方式来使用 CUDA\n",
    "- 加速 C/C++ 应用程序\n",
    "\n",
    "**PyCUDA**：\n",
    "- 完全对接 CUDA C/C++ API\n",
    "- 在 Python 中释放 NVIDIA GPU 性能的优先选项\n",
    "- 需在编写 C 代码，通常还要修改 Python 代码\n",
    "\n",
    "**Numba**：\n",
    "- 性能或许低于 PyCUDA\n",
    "- （当前）并未完全支持 CUDA C/C++ API\n",
    "- 仍能大幅加速应用程序，代码修改量小\n",
    "- 开发者可在 Python 框架下轻松编写代码\n",
    "- 能同时为 CPU 优化 Python 代码"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 第一步：编译CPU函数\n",
    "\n",
    "之前的课程中介绍过，Numba 可用于为 CPU 或 GPU 优化代码。在介绍 GPU 加速之前，我们首先编写第一个函数，然后为 **CPU** 编译该函数。如此一来，我们在学习 Numba 语法时便会更容易，而且稍后还可让我们有机会比较 CPU 优化的 Numba 代码与 GPU 加速的 Numba 代码的性能表现。\n",
    "\n",
    "通常，在对 Python 函数应用 [**函数装饰器**](https://en.wikipedia.org/wiki/Python_syntax_and_semantics#Decorators) 后，便可启用 Numba 编译器。装饰器即为函数修改器，使用十分简单的语法来转换所装饰的 Python 函数。我们将在此使用 Numba 的 CPU 编译装饰器 `@jit`："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 1,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "from numba import jit\n",
    "import math\n",
    "\n",
    "# This is the function decorator syntax and is equivalent to `hypot = jit(hypot)`.\n",
    "# The Numba compiler is just a function you can call whenever you want!\n",
    "@jit\n",
    "def hypot(x, y):\n",
    "    # Implementation from https://en.wikipedia.org/wiki/Hypot\n",
    "    x = abs(x);\n",
    "    y = abs(y);\n",
    "    t = min(x, y);\n",
    "    x = max(x, y);\n",
    "    t = t / x;\n",
    "    return x * math.sqrt(1+t*t)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "让我们尝试执行上面定义的求直角三角形斜边的运算："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "5.0"
      ]
     },
     "execution_count": 2,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "hypot(3.0, 4.0)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "下面我们将详细介绍调用 `hypot` 将会发生的具体过程，首先您要明确一点：初次调用 `hypot` 将会触发编译器，促使其编译函数代码。Numba 还会将函数的原始 Python 实现保存在 `.py_func` 属性中，以便我们能够调用原始 Python 代码，确保结果一致："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "5.0"
      ]
     },
     "execution_count": 3,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "hypot.py_func(3.0, 4.0)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 基准测试\n",
    "\n",
    "评测代码的性能是使用 Numba 的重要一环。下面让我们检测一下是否真正实现了加速。在 Jupyter Notebook（如本课程所用的笔记本）中执行此检测的最简单方法便是使用 [`%timeit`魔术函数](https://ipython.readthedocs.io/en/stable/interactive/magics.html#magic-timeit)。让我们先来测试一下原始 Python 的速度："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 4,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "674 ns ± 0.824 ns per loop (mean ± std. dev. of 7 runs, 1000000 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit hypot.py_func(3.0, 4.0)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "`%timeit` 魔术函数会多次运行该语句以准确估算运行时间。默认情况下，它会返回最佳时间，这有助于降低随机背景事件影响测量的几率。三选一的方法确保首次运行的编译时间不会被计算在程序运行时间内："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 5,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "195 ns ± 0.931 ns per loop (mean ± std. dev. of 7 runs, 10000000 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit hypot(3.0, 4.0)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Numba 对该函数加速效果良好，运行明显快于纯 Python 代码的实现。当然，`hypot` 函数已经存在于 Python 自带模块中，下面让我们进一步对比结果："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "128 ns ± 0.162 ns per loop (mean ± std. dev. of 7 runs, 10000000 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit math.hypot(3.0, 4.0)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Python 的内置函数甚至比 Numba 还快！这是因为 Numba 会为每次函数调用带来额外开销，这将超出 Python 自身函数调用的时间。即便是极快速的函数（如以上函数），也会受此影响。（附注：若您从一个 Numba 函数中调用另一个 Numba 函数，则调用开销会很小；若编译器将函数内联至另一个函数，调用开销甚至可忽略不计。简言之，请始终对您的函数执行基准测试，以检查加速效果。）"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 练习：使用 Numba 为 CPU 编译函数\n",
    "\n",
    "以下函数使用 [蒙特卡罗法确定圆周率](https://academo.org/demos/estimating-pi-monte-carlo/)（源代码来自 [Numba 主页](http://numba.pydata.org/)）。其正确性已有理论保证，无需担心具体的数学问题。\n",
    "\n",
    "请先完成两项 `TODO` 以使用 Numba 编译 `monte_carlo_pi`函数，然后再执行下方的 3 个代码单元，以完成以下任务：\n",
    "\n",
    "1. 确认编译版与未编译版运行结果一致。\n",
    "2. 对未编译版执行基准测试。\n",
    "3. 对编译版执行基准测试。\n",
    "\n",
    "如您遇到问题，请参阅 [此解决方案](../edit/solutions/monte_carlo_pi_solution.py)。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 7,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "nsamples = 1000000"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 8,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "# TODO: Import Numba's just-in-time compiler function\n",
    "import random\n",
    "from numba import jit\n",
    "\n",
    "\n",
    "# TODO: Use the Numba compiler to compile this function\n",
    "@jit\n",
    "def monte_carlo_pi(nsamples):\n",
    "    acc = 0\n",
    "    for i in range(nsamples):\n",
    "        x = random.random()\n",
    "        y = random.random()\n",
    "        if (x**2 + y**2) < 1.0:\n",
    "            acc += 1\n",
    "    return 4.0 * acc / nsamples"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 9,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "# We will use numpy's `testing` library to confirm compiled and uncompiled versions run the same\n",
    "from numpy import testing\n",
    "\n",
    "# This assertion will fail until you successfully complete the exercise one cell above\n",
    "testing.assert_almost_equal(monte_carlo_pi(nsamples), monte_carlo_pi.py_func(nsamples), decimal=2)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 10,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "10.9 ms ± 81.8 µs per loop (mean ± std. dev. of 7 runs, 100 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit monte_carlo_pi(nsamples)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 11,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "364 ms ± 5.91 ms per loop (mean ± std. dev. of 7 runs, 1 loop each)\n"
     ]
    }
   ],
   "source": [
    "%timeit monte_carlo_pi.py_func(nsamples)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Numba 工作原理\n",
    "\n",
    "现在，您已初步了解如何使用 Numba 编译器，下面我们来看一下其背后的工作原理。首次调用 Numba 包装的 `hypot` 函数时，将启动以下流程：\n",
    "\n",
    "![Numba 流程图](images/numba_flowchart.png \"编译流程\")\n",
    "\n",
    "使用 `.inspect_types()` 方法可得出类型推理结果，这将打印出源代码的注释版本："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 12,
   "metadata": {
    "scrolled": true
   },
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "hypot (float64, float64)\n",
      "--------------------------------------------------------------------------------\n",
      "# File: <ipython-input-1-1ae6c289554e>\n",
      "# --- LINE 6 --- \n",
      "# label 0\n",
      "\n",
      "@jit\n",
      "\n",
      "# --- LINE 7 --- \n",
      "\n",
      "def hypot(x, y):\n",
      "\n",
      "    # --- LINE 8 --- \n",
      "\n",
      "    # Implementation from https://en.wikipedia.org/wiki/Hypot\n",
      "\n",
      "    # --- LINE 9 --- \n",
      "    #   x = arg(0, name=x)  :: float64\n",
      "    #   y = arg(1, name=y)  :: float64\n",
      "    #   $0.1 = global(abs: <built-in function abs>)  :: Function(<built-in function abs>)\n",
      "    #   $0.3 = call $0.1(x, func=$0.1, args=[Var(x, <ipython-input-1-1ae6c289554e> (9))], kws=(), vararg=None)  :: (float64,) -> float64\n",
      "    #   del x\n",
      "    #   del $0.1\n",
      "    #   x.1 = $0.3  :: float64\n",
      "    #   del $0.3\n",
      "\n",
      "    x = abs(x);\n",
      "\n",
      "    # --- LINE 10 --- \n",
      "    #   $0.4 = global(abs: <built-in function abs>)  :: Function(<built-in function abs>)\n",
      "    #   $0.6 = call $0.4(y, func=$0.4, args=[Var(y, <ipython-input-1-1ae6c289554e> (9))], kws=(), vararg=None)  :: (float64,) -> float64\n",
      "    #   del y\n",
      "    #   del $0.4\n",
      "    #   y.1 = $0.6  :: float64\n",
      "    #   del $0.6\n",
      "\n",
      "    y = abs(y);\n",
      "\n",
      "    # --- LINE 11 --- \n",
      "    #   $0.7 = global(min: <built-in function min>)  :: Function(<built-in function min>)\n",
      "    #   $0.10 = call $0.7(x.1, y.1, func=$0.7, args=[Var(x.1, <ipython-input-1-1ae6c289554e> (9)), Var(y.1, <ipython-input-1-1ae6c289554e> (10))], kws=(), vararg=None)  :: (float64, float64) -> float64\n",
      "    #   del $0.7\n",
      "    #   t = $0.10  :: float64\n",
      "    #   del $0.10\n",
      "\n",
      "    t = min(x, y);\n",
      "\n",
      "    # --- LINE 12 --- \n",
      "    #   $0.11 = global(max: <built-in function max>)  :: Function(<built-in function max>)\n",
      "    #   $0.14 = call $0.11(x.1, y.1, func=$0.11, args=[Var(x.1, <ipython-input-1-1ae6c289554e> (9)), Var(y.1, <ipython-input-1-1ae6c289554e> (10))], kws=(), vararg=None)  :: (float64, float64) -> float64\n",
      "    #   del y.1\n",
      "    #   del x.1\n",
      "    #   del $0.11\n",
      "    #   x.2 = $0.14  :: float64\n",
      "    #   del $0.14\n",
      "\n",
      "    x = max(x, y);\n",
      "\n",
      "    # --- LINE 13 --- \n",
      "    #   $0.17 = t / x.2  :: float64\n",
      "    #   del t\n",
      "    #   t.1 = $0.17  :: float64\n",
      "    #   del $0.17\n",
      "\n",
      "    t = t / x;\n",
      "\n",
      "    # --- LINE 14 --- \n",
      "    #   $0.19 = global(math: <module 'math' from '/home/appuser/Miniconda3/lib/python3.6/lib-dynload/math.cpython-36m-x86_64-linux-gnu.so'>)  :: Module(<module 'math' from '/home/appuser/Miniconda3/lib/python3.6/lib-dynload/math.cpython-36m-x86_64-linux-gnu.so'>)\n",
      "    #   $0.20 = getattr(value=$0.19, attr=sqrt)  :: Function(<built-in function sqrt>)\n",
      "    #   del $0.19\n",
      "    #   $const0.21 = const(int, 1)  :: int64\n",
      "    #   $0.24 = t.1 * t.1  :: float64\n",
      "    #   del t.1\n",
      "    #   $0.25 = $const0.21 + $0.24  :: float64\n",
      "    #   del $const0.21\n",
      "    #   del $0.24\n",
      "    #   $0.26 = call $0.20($0.25, func=$0.20, args=[Var($0.25, <ipython-input-1-1ae6c289554e> (14))], kws=(), vararg=None)  :: (float64,) -> float64\n",
      "    #   del $0.25\n",
      "    #   del $0.20\n",
      "    #   $0.27 = x.2 * $0.26  :: float64\n",
      "    #   del x.2\n",
      "    #   del $0.26\n",
      "    #   $0.28 = cast(value=$0.27)  :: float64\n",
      "    #   del $0.27\n",
      "    #   return $0.28\n",
      "\n",
      "    return x * math.sqrt(1+t*t)\n",
      "\n",
      "\n",
      "================================================================================\n"
     ]
    }
   ],
   "source": [
    "hypot.inspect_types()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "注意，Numba 的类型名称往往与 [NumPy 类型名称](https://docs.scipy.org/doc/numpy-1.13.0/user/basics.types.html) 一致，因此 Python `float` 通常亦为 `float64`（又可称为“双精度”）。查看数据类型有时对于 GPU 代码相当重要，因为 `float32` 和 `float64` 的计算性能在 NVIDIA GPU 设备上相差甚远。若您的算法可以使用 `float32` 获取正确结果，那您就应使用该数据类型，因为转换至 `float64` 可能会显著降低函数速度。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## object 和 nopython 模式\n",
    "\n",
    "Numba 无法编译所有的 Python 代码。一方面，某些函数不具有 Numba 编译支持；另一方面，一些 Python 的数据类型（目前）无法高效编译。例如，Numba 不支持字典类型（至本课程编写时）。现在，我们尝试编译一串 Numba 尚无法编译的 Python 代码："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 13,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "'value'"
      ]
     },
     "execution_count": 13,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "@jit\n",
    "def cannot_compile(x):\n",
    "    return x['key']\n",
    "\n",
    "cannot_compile(dict(key='value'))"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "您可能很惊讶以上单元在执行期间未产生任何问题。这是因为，在默认情况下，Numba 会返回至 **object 模式**，而该模式不执行数据类型的专用化。object 模式旨在实现其他的 Numba 功能，但在许多情形下，您都需要从 Numba 获知它对数据类型的推理是否成功。您也可将 `nopython` 参数传递至装饰器，以强制进入 **nopython 模式**（另一种编译模式）："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 14,
   "metadata": {},
   "outputs": [
    {
     "ename": "TypingError",
     "evalue": "Failed in nopython mode pipeline (step: nopython frontend)\nInternal error at <numba.typeinfer.ArgConstraint object at 0x7f585e31e780>:\n--%<----------------------------------------------------------------------------\nTraceback (most recent call last):\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/errors.py\", line 599, in new_error_context\n    yield\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 199, in __call__\n    assert ty.is_precise()\nAssertionError\n\nDuring handling of the above exception, another exception occurred:\n\nTraceback (most recent call last):\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 142, in propagate\n    constraint(typeinfer)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 200, in __call__\n    typeinfer.add_type(self.dst, ty, loc=self.loc)\n  File \"/home/appuser/Miniconda3/lib/python3.6/contextlib.py\", line 99, in __exit__\n    self.gen.throw(type, value, traceback)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/errors.py\", line 607, in new_error_context\n    six.reraise(type(newerr), newerr, tb)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/six.py\", line 659, in reraise\n    raise value\nnumba.errors.InternalError: \n[1] During: typing of argument at <ipython-input-14-d3b98ca43e8a> (3)\n--%<----------------------------------------------------------------------------\n\n\nFile \"<ipython-input-14-d3b98ca43e8a>\", line 3:\ndef cannot_compile(x):\n    return x['key']\n    ^\n\nThis error may have been caused by the following argument(s):\n- argument 0: cannot determine Numba type of <class 'dict'>\n\nThis is not usually a problem with Numba itself but instead often caused by\nthe use of unsupported features or an issue in resolving types.\n\nTo see Python/NumPy features supported by the latest release of Numba visit:\nhttp://numba.pydata.org/numba-doc/dev/reference/pysupported.html\nand\nhttp://numba.pydata.org/numba-doc/dev/reference/numpysupported.html\n\nFor more information about typing errors and how to debug them visit:\nhttp://numba.pydata.org/numba-doc/latest/user/troubleshoot.html#my-code-doesn-t-compile\n\nIf you think your code should work with Numba, please report the error message\nand traceback, along with a minimal reproducer at:\nhttps://github.com/numba/numba/issues/new\n",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mTypingError\u001b[0m                               Traceback (most recent call last)",
      "\u001b[0;32m<ipython-input-14-d3b98ca43e8a>\u001b[0m in \u001b[0;36m<module>\u001b[0;34m()\u001b[0m\n\u001b[1;32m      3\u001b[0m     \u001b[0;32mreturn\u001b[0m \u001b[0mx\u001b[0m\u001b[0;34m[\u001b[0m\u001b[0;34m'key'\u001b[0m\u001b[0;34m]\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m      4\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m----> 5\u001b[0;31m \u001b[0mcannot_compile\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mdict\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mkey\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0;34m'value'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;32m/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/dispatcher.py\u001b[0m in \u001b[0;36m_compile_for_args\u001b[0;34m(self, *args, **kws)\u001b[0m\n\u001b[1;32m    346\u001b[0m                 \u001b[0me\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mpatch_message\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mmsg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    347\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 348\u001b[0;31m             \u001b[0merror_rewrite\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0me\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'typing'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    349\u001b[0m         \u001b[0;32mexcept\u001b[0m \u001b[0merrors\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mUnsupportedError\u001b[0m \u001b[0;32mas\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    350\u001b[0m             \u001b[0;31m# Something unsupported is present in the user code, add help info\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/dispatcher.py\u001b[0m in \u001b[0;36merror_rewrite\u001b[0;34m(e, issue_type)\u001b[0m\n\u001b[1;32m    313\u001b[0m                 \u001b[0;32mraise\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    314\u001b[0m             \u001b[0;32melse\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 315\u001b[0;31m                 \u001b[0mreraise\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mtype\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0me\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;32mNone\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    316\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    317\u001b[0m         \u001b[0margtypes\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;34m[\u001b[0m\u001b[0;34m]\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/six.py\u001b[0m in \u001b[0;36mreraise\u001b[0;34m(tp, value, tb)\u001b[0m\n\u001b[1;32m    656\u001b[0m             \u001b[0mvalue\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mtp\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    657\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0m__traceback__\u001b[0m \u001b[0;32mis\u001b[0m \u001b[0;32mnot\u001b[0m \u001b[0mtb\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 658\u001b[0;31m             \u001b[0;32mraise\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mwith_traceback\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mtb\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    659\u001b[0m         \u001b[0;32mraise\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    660\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;31mTypingError\u001b[0m: Failed in nopython mode pipeline (step: nopython frontend)\nInternal error at <numba.typeinfer.ArgConstraint object at 0x7f585e31e780>:\n--%<----------------------------------------------------------------------------\nTraceback (most recent call last):\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/errors.py\", line 599, in new_error_context\n    yield\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 199, in __call__\n    assert ty.is_precise()\nAssertionError\n\nDuring handling of the above exception, another exception occurred:\n\nTraceback (most recent call last):\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 142, in propagate\n    constraint(typeinfer)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 200, in __call__\n    typeinfer.add_type(self.dst, ty, loc=self.loc)\n  File \"/home/appuser/Miniconda3/lib/python3.6/contextlib.py\", line 99, in __exit__\n    self.gen.throw(type, value, traceback)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/errors.py\", line 607, in new_error_context\n    six.reraise(type(newerr), newerr, tb)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/six.py\", line 659, in reraise\n    raise value\nnumba.errors.InternalError: \n[1] During: typing of argument at <ipython-input-14-d3b98ca43e8a> (3)\n--%<----------------------------------------------------------------------------\n\n\nFile \"<ipython-input-14-d3b98ca43e8a>\", line 3:\ndef cannot_compile(x):\n    return x['key']\n    ^\n\nThis error may have been caused by the following argument(s):\n- argument 0: cannot determine Numba type of <class 'dict'>\n\nThis is not usually a problem with Numba itself but instead often caused by\nthe use of unsupported features or an issue in resolving types.\n\nTo see Python/NumPy features supported by the latest release of Numba visit:\nhttp://numba.pydata.org/numba-doc/dev/reference/pysupported.html\nand\nhttp://numba.pydata.org/numba-doc/dev/reference/numpysupported.html\n\nFor more information about typing errors and how to debug them visit:\nhttp://numba.pydata.org/numba-doc/latest/user/troubleshoot.html#my-code-doesn-t-compile\n\nIf you think your code should work with Numba, please report the error message\nand traceback, along with a minimal reproducer at:\nhttps://github.com/numba/numba/issues/new\n"
     ]
    }
   ],
   "source": [
    "@jit(nopython=True)\n",
    "def cannot_compile(x):\n",
    "    return x['key']\n",
    "\n",
    "cannot_compile(dict(key='value'))"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Numba 尝试编译函数时出现异常，在向下滚动至异常输出底部后，您将看到错误信息，其中会描述所发生的基本问题：\n",
    "```\n",
    "- argument 0: cannot determine Numba type of <class 'dict'>\n",
    "```\n",
    "\n",
    "**建议您选择 `nopython` 模式作为使用 `jit` 的最佳方法，因为这会为您带来更好的性能。**\n",
    "\n",
    "Numba 还会提供另一个装饰器 `njit`（`jit(nopython=True)` 的别名）："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 15,
   "metadata": {},
   "outputs": [
    {
     "ename": "TypingError",
     "evalue": "Failed in nopython mode pipeline (step: nopython frontend)\nInternal error at <numba.typeinfer.ArgConstraint object at 0x7f585e3c69b0>:\n--%<----------------------------------------------------------------------------\nTraceback (most recent call last):\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/errors.py\", line 599, in new_error_context\n    yield\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 199, in __call__\n    assert ty.is_precise()\nAssertionError\n\nDuring handling of the above exception, another exception occurred:\n\nTraceback (most recent call last):\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 142, in propagate\n    constraint(typeinfer)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 200, in __call__\n    typeinfer.add_type(self.dst, ty, loc=self.loc)\n  File \"/home/appuser/Miniconda3/lib/python3.6/contextlib.py\", line 99, in __exit__\n    self.gen.throw(type, value, traceback)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/errors.py\", line 607, in new_error_context\n    six.reraise(type(newerr), newerr, tb)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/six.py\", line 659, in reraise\n    raise value\nnumba.errors.InternalError: \n[1] During: typing of argument at <ipython-input-15-598d254e6e01> (5)\n--%<----------------------------------------------------------------------------\n\n\nFile \"<ipython-input-15-598d254e6e01>\", line 5:\ndef cannot_compile(x):\n    return x['key']\n    ^\n\nThis error may have been caused by the following argument(s):\n- argument 0: cannot determine Numba type of <class 'dict'>\n\nThis is not usually a problem with Numba itself but instead often caused by\nthe use of unsupported features or an issue in resolving types.\n\nTo see Python/NumPy features supported by the latest release of Numba visit:\nhttp://numba.pydata.org/numba-doc/dev/reference/pysupported.html\nand\nhttp://numba.pydata.org/numba-doc/dev/reference/numpysupported.html\n\nFor more information about typing errors and how to debug them visit:\nhttp://numba.pydata.org/numba-doc/latest/user/troubleshoot.html#my-code-doesn-t-compile\n\nIf you think your code should work with Numba, please report the error message\nand traceback, along with a minimal reproducer at:\nhttps://github.com/numba/numba/issues/new\n",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mTypingError\u001b[0m                               Traceback (most recent call last)",
      "\u001b[0;32m<ipython-input-15-598d254e6e01>\u001b[0m in \u001b[0;36m<module>\u001b[0;34m()\u001b[0m\n\u001b[1;32m      5\u001b[0m     \u001b[0;32mreturn\u001b[0m \u001b[0mx\u001b[0m\u001b[0;34m[\u001b[0m\u001b[0;34m'key'\u001b[0m\u001b[0;34m]\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m      6\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m----> 7\u001b[0;31m \u001b[0mcannot_compile\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mdict\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mkey\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0;34m'value'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;32m/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/dispatcher.py\u001b[0m in \u001b[0;36m_compile_for_args\u001b[0;34m(self, *args, **kws)\u001b[0m\n\u001b[1;32m    346\u001b[0m                 \u001b[0me\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mpatch_message\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mmsg\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    347\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 348\u001b[0;31m             \u001b[0merror_rewrite\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0me\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'typing'\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    349\u001b[0m         \u001b[0;32mexcept\u001b[0m \u001b[0merrors\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mUnsupportedError\u001b[0m \u001b[0;32mas\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    350\u001b[0m             \u001b[0;31m# Something unsupported is present in the user code, add help info\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/dispatcher.py\u001b[0m in \u001b[0;36merror_rewrite\u001b[0;34m(e, issue_type)\u001b[0m\n\u001b[1;32m    313\u001b[0m                 \u001b[0;32mraise\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    314\u001b[0m             \u001b[0;32melse\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 315\u001b[0;31m                 \u001b[0mreraise\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mtype\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0me\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0me\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;32mNone\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    316\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    317\u001b[0m         \u001b[0margtypes\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0;34m[\u001b[0m\u001b[0;34m]\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/six.py\u001b[0m in \u001b[0;36mreraise\u001b[0;34m(tp, value, tb)\u001b[0m\n\u001b[1;32m    656\u001b[0m             \u001b[0mvalue\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mtp\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    657\u001b[0m         \u001b[0;32mif\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0m__traceback__\u001b[0m \u001b[0;32mis\u001b[0m \u001b[0;32mnot\u001b[0m \u001b[0mtb\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 658\u001b[0;31m             \u001b[0;32mraise\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mwith_traceback\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mtb\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    659\u001b[0m         \u001b[0;32mraise\u001b[0m \u001b[0mvalue\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    660\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;31mTypingError\u001b[0m: Failed in nopython mode pipeline (step: nopython frontend)\nInternal error at <numba.typeinfer.ArgConstraint object at 0x7f585e3c69b0>:\n--%<----------------------------------------------------------------------------\nTraceback (most recent call last):\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/errors.py\", line 599, in new_error_context\n    yield\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 199, in __call__\n    assert ty.is_precise()\nAssertionError\n\nDuring handling of the above exception, another exception occurred:\n\nTraceback (most recent call last):\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 142, in propagate\n    constraint(typeinfer)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/typeinfer.py\", line 200, in __call__\n    typeinfer.add_type(self.dst, ty, loc=self.loc)\n  File \"/home/appuser/Miniconda3/lib/python3.6/contextlib.py\", line 99, in __exit__\n    self.gen.throw(type, value, traceback)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/errors.py\", line 607, in new_error_context\n    six.reraise(type(newerr), newerr, tb)\n  File \"/home/appuser/Miniconda3/lib/python3.6/site-packages/numba/six.py\", line 659, in reraise\n    raise value\nnumba.errors.InternalError: \n[1] During: typing of argument at <ipython-input-15-598d254e6e01> (5)\n--%<----------------------------------------------------------------------------\n\n\nFile \"<ipython-input-15-598d254e6e01>\", line 5:\ndef cannot_compile(x):\n    return x['key']\n    ^\n\nThis error may have been caused by the following argument(s):\n- argument 0: cannot determine Numba type of <class 'dict'>\n\nThis is not usually a problem with Numba itself but instead often caused by\nthe use of unsupported features or an issue in resolving types.\n\nTo see Python/NumPy features supported by the latest release of Numba visit:\nhttp://numba.pydata.org/numba-doc/dev/reference/pysupported.html\nand\nhttp://numba.pydata.org/numba-doc/dev/reference/numpysupported.html\n\nFor more information about typing errors and how to debug them visit:\nhttp://numba.pydata.org/numba-doc/latest/user/troubleshoot.html#my-code-doesn-t-compile\n\nIf you think your code should work with Numba, please report the error message\nand traceback, along with a minimal reproducer at:\nhttps://github.com/numba/numba/issues/new\n"
     ]
    }
   ],
   "source": [
    "from numba import njit\n",
    "\n",
    "@njit\n",
    "def cannot_compile(x):\n",
    "    return x['key']\n",
    "\n",
    "cannot_compile(dict(key='value'))"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "请参阅 [Numba 文件](https://numba.pydata.org/numba-doc/dev/reference/pysupported.html)，以获取 Numba 所支持的 Python 特性的详细信息。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 适用于 GPU 的 Numba 及 NumPy 通用函数 (ufunc) 简介\n",
    "\n",
    "下面我们会开始讲解 Numba 中的 GPU 编程，以及如何为 GPU 编译 [NumPy 通用函数 \\(ufunc\\)](https://docs.scipy.org/doc/numpy-1.15.1/reference/ufuncs.html)。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "开始学习之前，我们首先要明确 GPU 编程最重要的一点：GPU 硬件专为*数据并行*而设计。如要实现最大吞吐，GPU 需要同时对许多不同元素执行相同运算。\n",
    "\n",
    "NumPy 通用函数可对 NumPy 数组中的每个元素执行相同运算，本身就具备数据并行性，因此非常适于 GPU 编程。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 回顾 NumPy 通用函数 (ufunc) \n",
    "\n",
    "熟悉 NumPy 通用函数 (ufunc) 是本课程的先决条件。若您尚不熟悉该主题，或已遗忘所学内容，这里有一个非常简短的介绍。若阅读完本简介后，您感觉自己仍无法掌握对 NumPy 数组的使用和了解通用函数的基础机制，建议您抽出约 1 小时时间来学习 [NumPy 快速入门教程](https://docs.scipy.org/doc/numpy/user/quickstart.html)。\n",
    "\n",
    "NumPy 通用函数（“ufunc”）的概念是，此类函数可对相同或不同维度的 NumPy 数组或标量，进行逐元素的处理。\n",
    "\n",
    "作为示例，我们将使用 NumPy `add` ufunc 演示 ufunc 的基础机制："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 16,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([11, 22, 33, 44])"
      ]
     },
     "execution_count": 16,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "import numpy as np\n",
    "\n",
    "a = np.array([1, 2, 3, 4])\n",
    "b = np.array([10, 20, 30, 40])\n",
    "\n",
    "np.add(a, b) # Returns a new NumPy array resulting from adding every element in `a` to every element in `b`"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "ufunc 还可将标量与数组相结合："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 17,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([101, 102, 103, 104])"
      ]
     },
     "execution_count": 17,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "np.add(a, 100) # Returns a new NumPy array resulting from adding 100 to every element in `a`"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "对于维度不同但互相兼容的数组，我们还可通过名为 [*广播*](https://docs.scipy.org/doc/numpy-1.15.0/user/basics.broadcasting.html) 的技巧将其组合在一起。此方法会将低维数组复制多份，以匹配高维数组的维度。如有需要，请参阅 [`numpy.arange`](https://docs.scipy.org/doc/numpy-1.15.0/reference/generated/numpy.arange.html) 和 [`numpy.ndarray.reshape`](https://docs.scipy.org/doc/numpy-1.15.0/reference/generated/numpy.ndarray.reshape.html) 文档，本课程亦会多次用到这两种功能："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 18,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "c: [[ 0  1  2  3]\n",
      " [ 4  5  6  7]\n",
      " [ 8  9 10 11]\n",
      " [12 13 14 15]]\n"
     ]
    },
    {
     "data": {
      "text/plain": [
       "array([[10, 21, 32, 43],\n",
       "       [14, 25, 36, 47],\n",
       "       [18, 29, 40, 51],\n",
       "       [22, 33, 44, 55]])"
      ]
     },
     "execution_count": 18,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "c = np.arange(4*4).reshape((4,4))\n",
    "print('c:', c)\n",
    "\n",
    "np.add(b, c)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 为 GPU 创建 ufunc\n",
    "\n",
    "Numba 能够创建*经过编译的* ufunc，其过程涉及C代码，并不那么简单直观。借助 Numba，您只需实现一个对所有输入都执行的标量函数，并使用 `@vectorize` 装饰该函数。之后 Numba 便能确定出广播规则。若您熟悉 [NumPy 的 `vectorize`](https://docs.scipy.org/doc/numpy-1.15.0/reference/generated/numpy.vectorize.html)，那么肯定也能轻松掌握 Numba 的 `vectorize` 装饰器。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "在第一个示例中，我们将使用 `@vectorize` 装饰器为 **CPU** 编译和优化 ufunc。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 19,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "from numba import vectorize\n",
    "\n",
    "@vectorize\n",
    "def add_ten(num):\n",
    "    return num + 10 # This scalar operation will be performed on each element"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 20,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([10, 11, 12, 13, 14, 15, 16, 17, 18, 19])"
      ]
     },
     "execution_count": 20,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "nums = np.arange(10)\n",
    "add_ten(nums) # pass the whole array into the ufunc, it performs the operation on each element"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "我们将生成可在 GPU 上执行的 ufunc，它需要额外提供**显式类型签名**并设置 `target` 属性。类型签名用于描述ufuncs的参数和返回值使用的是哪种数据类型：\n",
    "```python\n",
    "'return_value_type(argument1_value_type, argument2_value_type, ...)'\n",
    "```\n",
    "\n",
    "请参阅 Numba 文档，获取有关 [可用的数据类型](https://numba.pydata.org/numba-doc/dev/reference/types.html) 及 [使用多个签名编写 ufunc](https://numba.pydata.org/numba-doc/dev/user/vectorize.html) 的更多信息。\n",
    "\n",
    "下面是一个简单的示例，展示了为支持 CUDA 的 GPU 设备而编译的 ufunc。该函数需要两个 `int64` 值做为参数，并会同样返回一个 `int64` 值："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 21,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "@vectorize(['int64(int64, int64)'], target='cuda') # Type signature and target are required for the GPU\n",
    "def add_ufunc(x, y):\n",
    "    return x + y"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 22,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([11, 22, 33, 44])"
      ]
     },
     "execution_count": 22,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "add_ufunc(a, b)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "虽然此次函数调用非常简单，但期间执行了多项操作！Numba 刚才已自动完成以下操作：\n",
    "\n",
    "* 已编译 CUDA 核函数，从而对所有输入元素并行执行 ufunc 运算。\n",
    "* 已为输入和输出分配 GPU 显存。\n",
    "* 已向 GPU 传入输入数据。\n",
    "* 根据给定的输入大小，已使用正确的核函数维度执行 CUDA 核函数（GPU 函数）。\n",
    "* 已将结果从 GPU 复制回 CPU。\n",
    "* 已将结果返回为主机上的 NumPy 数组。\n",
    " \n",
    "相较 C 语言的实现，以上 ufunc 的实现更为简洁。\n",
    "\n",
    "您可能想知道这一简单示例在 GPU 上的执行速度有多快。下面就让我们一探究竟："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 23,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "1.03 µs ± 0.728 ns per loop (mean ± std. dev. of 7 runs, 1000000 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit np.add(b, c)   # NumPy on CPU"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 24,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "689 µs ± 737 ns per loop (mean ± std. dev. of 7 runs, 1000 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit add_ufunc(b, c) # Numba on GPU"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "GPU 竟比 CPU *慢这么多*？？目前来说，这是正常现象。因为在本示例中，我们（故意）误用了 GPU，以帮助阐明哪些问题适用于 GPU 计算，而哪些最好由 CPU 执行：\n",
    "\n",
    "* **输入过小**：GPU 以并行方式实现高性能，可同时处理数千个值。我们的测试输入只分别采用了 4 和 16 个整数。只有大幅扩展数组规模，才能让 GPU 充分发挥作用。\n",
    "* **计算过于简单**：与调用 CPU 上的函数相比，到 GPU 上执行运算任务的额外开销很大。若计算过程不涉及大量数学运算（通常称为“计算密集型任务”），GPU 将耗费大部分时间等待数据传输完成。\n",
    "* **在 GPU 上来回复制数据**：某些情况下，对于单个函数而言，您可能值得耗费成本在 GPU 与 CPU 间来回复制数据，但我们更建议您对同一份数据运行多个 GPU 运算。针对此类情况，若将数据发送至 GPU，然后保留数据至所有计算结束，也不失为合理做法。\n",
    "* **数据类型大于实际所需**：本示例使用 `int64`，但我们无需这样做。使用 32 和 64 位数据类型的代码在 CPU 上的运行速度基本相同，整数类型在此方面的差异可能也不明显，但 64 位浮点数据类型却能为 GPU 带来巨大的性能消耗。相比 32 位浮点数，64 位浮点数的运算速度可能要慢 2 倍（基于 Pascal 架构的 Tesla）至 24 倍（基于 Maxwell 架构的 GeForce）不等。NumPy 在创建数组时默认使用 64 位数据类型，因此请务必按需设置 [`dtype`](https://docs.scipy.org/doc/numpy-1.14.0/reference/arrays.dtypes.html) 属性，或使用 [`ndarray.astype()`](https://docs.scipy.org/doc/numpy-1.15.0/reference/generated/numpy.ndarray.astype.html) 方法来指定使用 32 位数据类型。\n",
    "  \n",
    "  \n",
    "根据以上内容，我们来尝试一个在 GPU 上运行更快的示例。该示例将使用 32 位数据类型和更大的输入规模，并将以更高强度执行运算："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 25,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "import math # Note that for the CUDA target, we need to use the scalar functions from the math module, not NumPy\n",
    "\n",
    "SQRT_2PI = np.float32((2*math.pi)**0.5)  # Precompute this constant as a float32.  Numba will inline it at compile time.\n",
    "\n",
    "@vectorize(['float32(float32, float32, float32)'], target='cuda')\n",
    "def gaussian_pdf(x, mean, sigma):\n",
    "    '''Compute the value of a Gaussian probability density function at x with given mean and sigma.'''\n",
    "    return math.exp(-0.5 * ((x - mean) / sigma)**2) / (sigma * SQRT_2PI)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 26,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([0.11604577], dtype=float32)"
      ]
     },
     "execution_count": 26,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "import numpy as np\n",
    "# Evaluate the Gaussian a million times!\n",
    "x = np.random.uniform(-3, 3, size=1000000).astype(np.float32)\n",
    "mean = np.float32(0.0)\n",
    "sigma = np.float32(1.0)\n",
    "\n",
    "# Quick test on a single element just to make sure it works\n",
    "gaussian_pdf(x[0], 0.0, 1.0)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 27,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "66.4 ms ± 363 µs per loop (mean ± std. dev. of 7 runs, 10 loops each)\n"
     ]
    }
   ],
   "source": [
    "import scipy.stats # for definition of gaussian distribution, so we can compare CPU to GPU time\n",
    "norm_pdf = scipy.stats.norm\n",
    "%timeit norm_pdf.pdf(x, loc=mean, scale=sigma)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 28,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "4.23 ms ± 92.5 µs per loop (mean ± std. dev. of 7 runs, 100 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit gaussian_pdf(x, mean, sigma)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "尽管在 GPU 上来回复制数据时增加了额外开销，但此示例依旧取得了显著的速度提升。对大型数据集使用特殊函数（`exp`、`sin` 和 `cos` 等）的 ufunc 尤其适合在 GPU 上运行。\n",
    "\n",
    "做为最后的对比，让我们就 Numba 专为 **CPU** 优化的 `gaussian_pdf` 函数作出定义并计时："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 29,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "@vectorize\n",
    "def cpu_gaussian_pdf(x, mean, sigma):\n",
    "    '''Compute the value of a Gaussian probability density function at x with given mean and sigma.'''\n",
    "    return math.exp(-0.5 * ((x - mean) / sigma)**2) / (sigma * SQRT_2PI)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 30,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "23.1 ms ± 24.9 µs per loop (mean ± std. dev. of 7 runs, 10 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit cpu_gaussian_pdf(x, mean, sigma)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "以上运行速度大大优于未编译的 CPU 版本，但又远低于 GPU 加速版本。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## CUDA 设备函数\n",
    "\n",
    "ufunc 非常适合执行元素级的运算，这一任务也十分普遍。然而，许多函数并非如此。如要为 GPU 编译的函数**并非是**元素级的向量化函数，我们可以使用 `numba.cuda.jit`。在本课程下一节内容中，我们将深入学习 `numba.cuda.jit`，但现在让我们演示一下如何使用它来装饰辅助函数，以供GPU加速的ufunc使用，这样您就不需要将所有逻辑都塞入单个ufunc的定义中了。\n",
    "\n",
    "注意，下方的 `polar_to_cartesian` 无需使用类型签名，此外它还传递两个标量值，这与我们所使用的向量化的 ufunc（类似于下方的 `polar_distance`）不同，因为后者需将 NumPy 数组用作参数。\n",
    "\n",
    "`device=True` 参数表示装饰后的函数**仅**可从在 GPU 上运行的函数所调用，而无法从 CPU 主机代码调用："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 31,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "from numba import cuda\n",
    "\n",
    "@cuda.jit(device=True)\n",
    "def polar_to_cartesian(rho, theta):\n",
    "    x = rho * math.cos(theta)\n",
    "    y = rho * math.sin(theta)\n",
    "    return x, y\n",
    "\n",
    "@vectorize(['float32(float32, float32, float32, float32)'], target='cuda')\n",
    "def polar_distance(rho1, theta1, rho2, theta2):\n",
    "    x1, y1 = polar_to_cartesian(rho1, theta1) # We can use device functions inside our GPU ufuncs\n",
    "    x2, y2 = polar_to_cartesian(rho2, theta2)\n",
    "    \n",
    "    return ((x1 - x2)**2 + (y1 - y2)**2)**0.5"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 32,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "n = 1000000\n",
    "rho1 = np.random.uniform(0.5, 1.5, size=n).astype(np.float32)\n",
    "theta1 = np.random.uniform(-np.pi, np.pi, size=n).astype(np.float32)\n",
    "rho2 = np.random.uniform(0.5, 1.5, size=n).astype(np.float32)\n",
    "theta2 = np.random.uniform(-np.pi, np.pi, size=n).astype(np.float32)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 33,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([1.2485787 , 0.7118244 , 1.4680566 , ..., 1.5220381 , 1.3889159 ,\n",
       "       0.69207233], dtype=float32)"
      ]
     },
     "execution_count": 33,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "polar_distance(rho1, theta1, rho2, theta2)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "请注意，CUDA 编译器将主动内联(inline)设备函数，因此函数调用通常不会产生任何额外开销。同样，`polar_to_cartesian` 返回的“元组”（tuple类型）实际上并不是作为 Python 对象而创建的，而是临时表示为结构体（struct类型），然后由编译器对其进行优化。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## GPU 所支持的 Python\n",
    "\n",
    "Numba 在 CPU 上受到限制无法编译所有函数，GPU 上同样如此。GPU 所支持的 Python 操作包括：\n",
    "\n",
    "* `if`/`elif`/`else`\n",
    "* `while` 和 `for` 循环\n",
    "* 基础数学运算符\n",
    "* 选自 `math` 和 `cmath` 模块的函数\n",
    "* 元组\n",
    "\n",
    "详情请参阅 [Numba 手册](http://numba.pydata.org/numba-doc/latest/cuda/cudapysupported.html)。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 练习：使用 GPU 加速函数\n",
    "\n",
    "我们来使用 GPU 加速“零抑制”函数。这是处理波形的一个常见运算：将所有低于某绝对量的样本值强制设为零，从而消除低振幅噪音。下面就让我们创建一些样本数据："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 34,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "[<matplotlib.lines.Line2D at 0x7f582c1023c8>]"
      ]
     },
     "execution_count": 34,
     "metadata": {},
     "output_type": "execute_result"
    },
    {
     "data": {
      "image/png": "iVBORw0KGgoAAAANSUhEUgAAAXoAAAD8CAYAAAB5Pm/hAAAABHNCSVQICAgIfAhkiAAAAAlwSFlz\nAAALEgAACxIB0t1+/AAAADl0RVh0U29mdHdhcmUAbWF0cGxvdGxpYiB2ZXJzaW9uIDIuMS4yLCBo\ndHRwOi8vbWF0cGxvdGxpYi5vcmcvNQv5yAAAIABJREFUeJzt3Xl8VNX5P/DPA0lYwhoIayAh7MhO\nRBRFRVyAKq5Va92t2qqt2p8Fa22rVstXW2ut+45116ogoKCA7AJhCzsECCRsSYCEAIEsc35/5Gbm\nTDKT2e69Z+65z/v1yit37tyZ+8z2zJ1zz3kOCSHAGGNMX41UB8AYY8xanOgZY0xznOgZY0xznOgZ\nY0xznOgZY0xznOgZY0xznOgZY0xznOgZY0xznOgZY0xzCaoDAID27duLjIwM1WEwxpijrF69ulgI\nkRpqu7hI9BkZGcjOzlYdBmOMOQoR7QlnO266YYwxzXGiZ4wxzXGiZ4wxzXGiZ4wxzXGiZ4wxzXGi\nZ4wxzXGiZ4wxzXGiZ4y50oJthSg4elJ1GLbgRC/JKSjBJyv3qg6DMWaD299dhcteWKw6DFvExcjY\neHHFS0sBAMt3HcYvRnbHWZntFEfkXIVlp/Dwp+vxnxuHoW1ykupwHK2iyoNJLy/FjkNlGJTWGm/f\neiZS+DmNyYvzdgAAjp+uUhyJPfiIPoDp6/bjtndXqQ7D0Z6ZtQVLcosx7Knv0eexb1WH42gb9pVi\ny4FjqPIIrN1bgo9WhDXqnQWxaX8pnv9+u+owbMWJPojyympkPjpLdRiO9fW6/d7limoPLnhugcJo\nnMvjEdhfUu63rryyWlE0znew9BQmvrjEb92CrYWKorEPJ3rDstzieus8QkEgGig7VVlvXd5hd5z0\nMtu/ftiOBz5e67fu01X5iqJxvlF/n1dv3f/7fL2CSOzFid4wc8MB1SFo47rXlgdc7+Fvzoj9Z35u\nvXXFxyswm9+vEftqbUHA9YdPVNgcif040Rs+WhG4t817S3fbHImzlVdUY+vBsoDXnRXgaIpF5zcf\nrlEdguM89Kn+R+7BcKIHUHqyflNDrb9+s9nGSJzvtYU7g15XVHbaxkgYC1/hsVOqQ7AUJ3oAN7+z\nQnUI2liz96jqELTx/NxtqkNwjXeX5akOwVKc6AHkFJSqDkEbuYXHVYegjRcDtM/LSsuD/xJl/n7Y\nfKjB6/ce0buzgOsTvRChTxDmFZ+wIRI9HCht+Cfws99ttSkS/Q15Yi5OV3FXy3B8sTrwidhas3L0\nPrnt+kR/+UtLQm7z+PSNNkTifOF8ab7y486wtmPh+fts/uIMx3ebDqoOQSnXJ/qN+46F3Gbxjvp9\n7Fl91WF2n+Q8b55d/GvTNKv3HFEdgmVcn+iZearDzOAezvQhhTvmYNH2IosjcY9rXg08/kMHIRM9\nEXUjogVEtIWINhHR74z1KUT0PRHtMP63NdYTEb1IRLlElENEw61+ECw+/LbOCM5gCo6Wh97I5eZu\ndndTg5m4qTC8I/oqAL8XQvQHMArAfUQ0AMAUAPOEEL0BzDMuA8B4AL2Nv7sBvGp61CwuzdnUcM+G\nWhf840drA9HAvR/wgCizXPrCItUhKBcy0QshDggh1hjLZQC2AOgKYBKAacZm0wBcaSxPAvC+qPET\ngDZE1Nn0yG323Ua9z8ozpqvth8Lv8qvr0X9EbfRElAFgGIAVADoKIQ4ANV8GADoYm3UFIFddKjDW\n1b2vu4kom4iyi4rip50x56+XBFzPR1hMtRV/vEh1CNogCrx+pqbdLMNO9ETUAsD/ADwohGioq0qg\np7De16QQ4g0hRJYQIis1NTXcMCzXqmki8qZOVB2GFn54+Hx+Lk3SpnkiOrZqqjoMbez++8SAX5zF\nx/Us0xFWoieiRNQk+Q+FEF8aqw/VNskY/2uLOhcA6CbdPA3AfjCtBZqpJ7N9MgBgYNdWdofjaCUn\n61dTHNu35gfztDtG1rtO1+RklZkPnAsAAb84/23MPKWbcHrdEIC3AWwRQjwvXTUDwK3G8q0Apkvr\nbzF634wCUFrbxBNv8hsY9tynY4t669wy7Vg0As2126hRzY+7b+4/1+5wHO2+j+o3Ez591SAAwPl9\n6v/6XcLjPCIysGvroNeVNFDg0MnCOaIfDeBmAGOJaJ3xNwHAVAAXE9EOABcblwFgNoBdAHIBvAng\nN+aHbY7znvWf9WjG/aO9y40CNOKt3sMFu4L526wtQa8jItx2TobfulM8S1JQS3MP11vXLKlx0O25\nDEJwDVWmBYArh3axKRK1wul1s0QIQUKIwUKIocbfbCHEYSHERUKI3sb/I8b2QghxnxCipxBikBAi\n2/qHYY7BaW28y1cOq3f+GAVH9S58ZKYurf1/Ft8xuoff5Zvf5oqh4bp6uP978cf/d4Hf5cn/22Bj\nNM7yQYj5dZ+7bohNkajl2pGxxwJMdye7Z0xmvXWa9ryyxMI/XOh3uX3LJL/Lq/L411G4nrvWPxl1\nbsMnZcP13JyGSz0nNnZHCnTHowwgSO8q3/VE6Neppd+6tXtLrAtIM3U/QM2TEhRF4nyNG/m/W5sk\nBG/GYQ3LTE1WHYISrk30dT88gbxz25l+l/+3puFSp25VUeVRHQJjYfn8nrNVh6CEaxN9dp2mg0cu\n7Vtvmy5tmtkVjqPV7d5Xt0251qjMFL/L5RV8EjFa4/p38LvcUA8y5tOuRZN6676+b7Tf5QOl+tVi\ncm2iv+WdlX6XExsHPsJ/9SauyRZK3fLEgboAAsAfLuvndznUeRI3Wrnbv1RuUkLgj+gDY3v7Xd60\nP3S5bbfr2Kp+kgeAod3a+F3eFmRyeydzbaKv66we7QKuHz/I8WV6LFdep6vkpKGBj+iHpvl/oLhc\ncX1bD4aXsIfUSU4BBp+7Xt0DkCuGhNeV8h8aztXLid5Q/4MTGHexrO/3n60Pa7tGdc6LcJ6v78gJ\n/1Gxf7l8QFi3W5fP8x7X9WGdrpVjgvzSrCucyYichhM9grcpB/J0AwOD3GrDvuiSzOHj9Yf6u90L\nP/gPwb/prPSwbvfawp1WhONof56+ye9yQiP3pjv3PnJJVnpK6I0MldXcw6Qh7QOc7Arm/3iicGaj\ns3oE/5w/e+1gGyOxnysT/dE6P48HpwWvfVHXvpJTZoejlYWPXBD2tktyuUYLs0/dpkPZtcPTbIzE\nfq5M9AeP+Sfr1JbhH4VuOaBf+52ZkpvwwCizTL16kOoQXKOhLwEduDLRV1X7nwUMVec7UGlYFp0H\nxvZSHYJjXDawU4PXL50y1qZImNO5MtF/8FPDhY7qqtstkPlE2hc+JTkp9EYu5anTHbBN84afq648\noC9sf5zQL/RGGnNloo+0rnzr5okWReJ8L83PjWj7a0fo3RYai4Xb42dKTd1E0jwL6FdG25WJfulO\n30nA4d35aD0WbyzaFdH2LZv6f2nuKgp/4mbdfbKq/uQtkeAeYcFdPjiyuvM7IphQ3AlcmejlWWTe\nvCVLYSTs3g9Wqw4hbsSaXDZzGQSvus1gCRGWI77+jeVmhqOcKxO9rGli5CVf84pPWBCJ853TM3AZ\niYZs1+zIKRa7onhfje7le86nfMkTkNTKiWIQnzxy9qRmBfdcl+jrTi3WJEjRqIZs1bDokRkamouT\nRabu1IvBjB/oq8XEXX99Xv0xsnNHAPDUpDMsiCQ+uC7RV9Rpx4z0Jx0A/PaTtWaFo5VLz+gY1nbh\nFpdys0vCfC6vCjDlJQPmbDrkXR4S5oDIukXQdOK6RB9gzu+I8UQbgQ3r1jas7Uakh7edm43MCK8s\nBw9QC21UmE2KaW2bWxyJOq5L9Mw64Y4uvHFkd4sjcb5ofmmywML90gxW+18H+j6yIKKtgf77i/uY\nHIl76fyBYvFnbL8OoTfSnOs+cb+alh3V7e44t4fJkThfbiGflGbxj6JsrxUaTZjgukS/viC62unc\nFlpf/hH95tZ0uroTl7Dofb66QHUIpnFdomfm+fOMjabcj869HcJ1usqcftu3vbsy9EYsKHkSoj2H\n9Rkv4+pE//m9Z6sOwdHMOqLftJ+nwfvCpKPHnYU8AO1kRWS1rGT3X+irrvryAn1m7XJ1oj8zzLPx\nLLRYTlZn5x01MRJneuwr36+jy85ouDxxQ05oNqIzGnsORz+vs66/LV2d6GORf4QnCZc9cFHviLaf\nfJmvbGysxbx0c9d5kZ34/12Ez73uGpkxWEYznOijdOgYTykYi7Myfb+muN6Nv24pkQ3ciaZek85O\nSE03OX+9JKLbRjLnsZNwoo/SXj6iZxYJNeNZXRMGRd/Uo6OrX1nmXY60h2TrZnrOPeGqRB/rZAI7\nn5ngXX74s/WxhuNq/PPaPOntklWHELdacrdoAC5L9Gv2xnbSr7HmEwjbKdxCU4zFQvdJv8PlqkTP\nxcjiR7SjFRljkQuZ6InoHSIqJKKN0rq/EtE+Ilpn/E2QrnuUiHKJaBsRXWpV4NE4VcmJnjHmPuEc\n0b8H4LIA6/8lhBhq/M0GACIaAOAGAGcYt3mFiOKmS8Bbi33zm/7nxmEKI3E+np/UPHWnvWPx47uN\nB1SHYIqQiV4IsQjAkTDvbxKAT4QQp4UQuwHkAhgZQ3ymyt7ja6Mf3au9wkicT+5e+usLeiqMxPm+\n23Qw5vtISU7yLvMXR2y6tmnmXb73gzUKIzFPLG309xNRjtG0UzuTRFcA+dI2Bca6uCN/MFjkzv2/\nBd7lq02Y5ajcxSM650qJPtoZo87O9E2usYknCY/JizcOVR2C6aJN9K8C6AlgKIADAP5prA90hi3g\n4QUR3U1E2USUXVRUFGUYLB6YUdkzp6DEhEicqUo6Ar+of3S101s1870GVR5uVovF8O76zYAWVaIX\nQhwSQlQLITwA3oSveaYAQDdp0zQA+4PcxxtCiCwhRFZqamqgTZhDdJF+6karysXNDTNzfO3AozLD\nm/auLnkavG/W69GuHA25C3XHVtGNctWxR1hUiZ6IOksXrwJQ2yNnBoAbiKgJEfUA0BuAVnVTr8/y\nfY8VHz+tMBLnu/d8X9v++8vzlMURT6Idgi8PQPssO7+BLfU2S/rSvOmsdIWRxJdwuld+DGA5gL5E\nVEBEdwJ4log2EFEOgAsBPAQAQohNAD4DsBnAdwDuE0Jo1fjaNNH3lH27wb1HTmYY2q2Nd3njPm5X\njsUFfX2/io+fjr5Mr9O9vWS3d/nmUZzoa4VsXBVC3Bhg9dsNbP80gKdjCSqeTRrWFdOW7wEA7CzS\nZ2ICFQZLo2P3lfBsVbHo37mV6hDiTlvucOHlmpGxsda5qTVMOgp9b1meKffpNGbNCNW5dWTFuxhj\n0XFNos8zaVowHU/URCqWGXxk/FwyZg/XJPrPVukz0a9q7u0fw5gzuSbRL9tZ7F2We3uwyMk1vv92\n5UB1gTBmg70xTE0YL1yT6LceLPMun9+H++3HQkiZfkS6foNLGJN9mu38qS5dk+hlZ2ZwcopF/hHu\nIROPzunpG2x1mMd4xOTDu87yLr+8YKfCSMzhykSf0NiVD9s0Mzf4Bjv3TG2hMBIma57kKxRbcJS/\njGOhW9FDzngsYq8v3BV6IxaWn3Yd9i5/JB1FRmP8QN+A9Y9XOr+5gZmHEz2LSawztQ3q6hs0tTS3\nuIEt9TR9ne/XUYcIJwWvKzHB93HefqisgS31tEL60pSbsRgn+qh0kj6QZg0ecqpYm8F6pvomtr79\n3VWxhuM48kC+pBify0TpW3fNXvdVA5Xnm/jvnbH9OtINJ/ooTL1mkHdZ7oHCIvfzM31F4ipcOGtV\n2Snf4LMOUVZbrOX2d2L+EV83yMY8KbgfTvRR6CQN3XdzASkzyIXN3Gjh9kLvctPE2GbdHOPybsOf\nrHJv1c5QXJHozapzU0ue12HBtsLgG7KQEhq54i0YVGW1ecfhLUyYAIbpyRWfssMnKky9P4/UXLNa\nahdkkUtKcMVbkDGlXPEpa2xy8Sx5cog5mw6Zet+MMWY2VyT6dfnmHnXLbfRFZTwCkTEW31yR6I+c\nqPQut26WqDASxphTjB/Yybu89aCzZ0BzRaLfK3W7um5EmsJInE+uRd+OZ/BhGrt6uC9XOH28jCsS\n/dtLfEP2J4/vpzAS58spKPUuZ0qDnVhsHrm0r+n3WV6h1XTNDfJYkIjlXkyJDq+P5ezowyS/B5z+\ngqkmD0SZdsdIhZE437FTvibFgVIpiFj8dmwv7/LhE+45fyQPtkttGdvAM9/9+H6xJjh8AJYrsp7T\nf3bFk0ZSD6bmSeb3266ocs/o2MPHfd1+h6SZk+h/N66Pd7mRi6ZqlLs83zMm05T77NWhpXc5O8/Z\n3ahdkeiZeXYVHTf9Pn9xVnfvcqWLyiBUWfBY5V9c5SYPFIxn8sGcFXMRT/4yx/T7tBMn+ijJQ/et\naB+MV498Yf4b/iYp0XtcVDtIrlxJMD85Xf/6T6bfZ7zKLfQdgFjRyuL0tyUn+ihNkU7qyjXFWeTk\nAWgnXXQC8c3F1tb1L3bRLFNyKYmJgzo3sKU7caKPUjOpANW8rVzvJhYdpbLP2w66p476aRedj7Ca\nPJdBrHX9dcSJPkryiS43NTdYLafAfXXUAcCClhtX+fe8HapDiGuc6KPUuY3vqMFNPUWstnB7keoQ\nlGiayB9FZh1+d0VJblf+cAXPz2mWVQ7vxhatJgmx1aJnrCGc6BljTHOuSvQvXD9UdQiMMQcZ17+j\n6hBMoX2iLy33DTPv0qaZwkj0MrJHiuoQGLNcRrvm3uWSk+ZOYGQn7RP9419v9C7zRN7muWN0huoQ\nGLPcNVK12wOlpxRGEhvtE/2M9b7Rh5zmYyMX4erXqZXCSPRyxZAupt7fmRltTb0/N5OLmTV2cGEz\n7RO9jA/oY3NKGrWaLv2kZZGTa/qYPQ7jrvN8Rb3cUDvIynLM1dJr4+TiiCETPRG9Q0SFRLRRWpdC\nRN8T0Q7jf1tjPRHRi0SUS0Q5RDTcyuAjNax7m9AbsaB2SPVEzC4cdX1WN1PvL97JCfiqYV1Nve+x\n/Tp4l3cXnzD1vuNRlce6LzO5BtFcB88PHc4R/XsALquzbgqAeUKI3gDmGZcBYDyA3sbf3QBeNSdM\nczRN5L7KsbjprRWW3fdZme46ubtkh2/IfkeTh+zLcy7sLDS/2mi8kY+zn7lqkKn33adjC+9y/tGT\nDWwZ30ImeiHEIgBH6qyeBGCasTwNwJXS+vdFjZ8AtCEirjDEQnJy+2c0HpM6CZg16UggczYdtOy+\n44Xc8nX5EHPTjfzL9YvVBabet52ibaPvKIQ4AADG/9rfil0B5EvbFRjr6iGiu4kom4iyi4qcOez9\nTxP7qw5BG72lSR7coKjMnsqSVQ5uV45GUoKrTjuGzexnJdBhWcB3mhDiDSFElhAiKzU11eQw7DEq\ns53qELQxoIuvF8+WA8cURqKXslNVoTdyOLnbNJeSCCzaRH+otknG+F9bp7cAgHxWLQ3AfmhKbm7g\nPvrmOejg/srxxg1F4l5ekKs6hLgXbaKfAeBWY/lWANOl9bcYvW9GASitbeLRUe8OvhM1JScrG9iS\nReL291apDoE5yJuLd6sOIe6FnN2ZiD4GcAGA9kRUAOAvAKYC+IyI7gSwF8B1xuazAUwAkAvgJIDb\nLYg5biQ05vZAxlj8C5nohRA3BrnqogDbCgD3xRqUFe42aWb4YF74YTuemDTQ0n0wxlg0XHNIOjjN\nui5sADBt+R5L7z+eNE/iE16MOYnWid4jdS0b3t3a+h9NXNSti7uwMTfRYYiH1p/YCmmYuckj9utx\n00TPCY20ftvYqkf7ZNUhsBCeu3aI6hBipvUntvCYb1CK20ZeWmlsP2eOe4gXcldcq5oUk7ijgGku\nHdhJdQgx0/rdMOa5Bd7lDi3NrSfiZmltuXJlLOTBqi2bhuwPEZXHf8ajts0ilyo+XWVdpUwraZ3o\nmXn2l5R7l285O92SfVjdvBYv5PIH5/Rsb8k+uICfeRpJb8xj5c4cacyJnoXl3g9We5fbNE+yZB8P\nXNjLkvuNN1+u9RXHOr+PNc1gbqtxYyW52depI4050cfocpNnB4pXOQWllu/jkjN8baHyXL+62V3k\nqxGf3MSappv2LZpYcr/xrH0Law5A5NN7G/dZ/zmwAif6GP1iZHfVIWijl1RSwqPxEakdR9vn9PQV\n3KtywSxTAPC3K82tRV9LLlX83rI8S/ZhNU70MUpr20x1CNqQ25XLK5150iscG2w4KpR/KXy/2bkz\nI4Vy4rSvzdyqZjAdcKKPUSPutmmJT1buVR2CZXJtnvXp1x+usXV/dpKnSnTLyfxocKKPkdw9ruRk\nhcJI9OLUk17MXgVHfb3BnDx5t9U40ceoVdNE7zKXKjYPf2RZOPZJ3X4TGvMhfTCc6E3020/Wqg5B\nGx6eyIWFQe4Fw7NLBeeKRN/Soi5sddnRBdEt9hSfVB0Cc4Cv1u5THYIjuCLRW9VXmVmn7LQzRyAy\nFo9ckejvOd/aSUcYY3qzqiaRXVyR6G87J0N1CIwxB7u4f0fVIcRE20Qvl4Il7mBrmp8N7qw6BMZs\nd5GU6J04alvbRL9x3zHVIWgpkyfKiMnJCj734HROLBinbaLPP8q9NqzQ1OL5Ypu5qLwu/9B0DrmP\nvhO7/mqb6H+j8bBvlVo3Swy9UQz6dmpp6f2rJtc2t3qKumevGWzp/bvJOKnpRp65zim0TfTMPLNy\nDniXbzjT2mqdrSz+IlFNPhjs2MraUsIZLmpmu9WiyXBqyTXpP812Xh0mTvQspK+kiTKsnnv3rnN7\neJfX7j1q6b5Us3qS9T4dfWWfda7vDwB3nmtfF+qXF+y0bV9m4URvgj9O6Kc6BEtVVtvXJpkoTWqd\nLxWs0oVcoliuv2+FZtL5lBnr91u6LxUWSYXvuqVwufCGcKI3QTfNJ8s+dOyUbfsalNbauywceNIr\nlIPSc5na0tqmG7n2y+Nfb7R0XyoclarFchfqhnGiN0GTRL2fxq0Hy2zbVwvNy1V8taYg9EYsLBVV\n7pg5ywx6ZyibtG7mm6vyQKl+zQ2qFJU5r3dDKAu2cZ19s3D9+fBxojfBiPS23uWX5ucqjEQvTp2I\nmdnjhy2FqkNwDE70Jluzt0R1CNro0d7ak5XM2bL3HFEdgmNwojcZnxIyT+c2TVWHwOIYN92ET/tE\nP0TqxWEHi7tGu8qz321THQKLY2WnuG5QuLRPS40sHuBTFxdTM0/xcf1OxjI9VFU7q8dPTImeiPKI\naAMRrSOibGNdChF9T0Q7jP9tQ92Plbq05oEUZmnTXO/yBIyF60RFteoQImLGEf2FQoihQogs4/IU\nAPOEEL0BzDMuK/PLUdbWwHCT5CS9+7gz1pD2LXzdqK0uBWI2K5puJgGYZixPA3ClBfto0Oo9vhop\nTp8CLJ4M6NJKdQiMKTOgi+98X2OHjcSNNdELAHOJaDUR3W2s6yiEOAAAxv8OMe4jYrmFvpGcDns9\n4ppcJIs5w8UDnD0FXjz5xchu3uV9Jc4aGBlroh8thBgOYDyA+4hoTLg3JKK7iSibiLKLiswdLfj6\nol3eZasLR+nu6AlfPZFbbZp7N6mx9n0E8POsNFv28+h4vQvu2WlEeop3+XSVi9rohRD7jf+FAL4C\nMBLAISLqDADG/4DD14QQbwghsoQQWampqbGEUc+uohPeZbmwE4vc3iO+mbo6tLSnX7uuv8LknhoP\njO1tyz6tLpwWDwZ1tacLtdwu77R6e1EneiJKJqKWtcsALgGwEcAMALcam90KYHqsQTJ1Jr281PZ9\ndm6t50CpQ1LtnqY2TZlo135Ukk+SWqlJgi9dvrcsz5Z9miWWI/qOAJYQ0XoAKwHMEkJ8B2AqgIuJ\naAeAi43LjIVtdK/2qkOwxGFpXIBd844muqAZrF9nezoJJEuVVb9Y7awqpFF3SRFC7AJQb9JLIcRh\nABfFEpQTdWjZBIUaVltk5pFzu4oJpqs9wnHdAsNxfh9zm351pP/XvU24v755HNb8GbZqKbmnJNvT\n3CB7auZm2/dph/42HdE7GSd6k9jVI8UNHhznO1F5qtJZvRsaslaqbKqik4BOcyXIs4+11nxCeTNw\nojdJM+mkl07JSQW5d49ONelVH1HP2XRI6f7NdJpnl4oIJ3qTyG2f2w/ZN/We7nYVnwi9EXOdpbnF\nqkNwFK0TfWJj+048yee49pfYN5m27nQ6omfm+ZGnZIyI1om+stq+03ryLPTf5Oy3bb+6e3/5HtUh\nsDjktBIEqmmd6Id1b6Nkv/v5TciYpeZv5fliI6F1or/hzG6hN7LA6Uo+UcSYjpxaUkLrRH9ebzUD\nKZxW8IgxFjm54F+80zrRq7KzSL+eIkO7qWkGYyye9O3Y0rucd9g5n3PtEn2+VG2RJx0xz1OTBqoO\nQRv8vnSua0fYU17abNol+pKTld7llk15xJxZurbluXedir9YzOPUZlntEr3QtlKK/eQ3dbWHn1ez\nvHjjMFv359QTiPGoQuqy7aRPhHaJ/kApD1Yyy/ebfUPmdZ0MRIUBNhfheuPmLFv3Z6dzerazdX+T\nhnbxLjtp8hHtEv1fpm9SHYI25KP49i3sPSq8e0ymrfuz07HyytAbmUjn6TTtrlzZXKpptavouK37\njoV2if7gMT6iN4uKmum10ts1V7ZvK5w4XeVdbt5EXZu5bgX3urSx99xRgjSRS3beUVv3HQvtEr1K\nr940XHUIplI58KtXql5Hoct3HvYud7U5Ocmc1NwQjn6dWobeyCJVDjpvxYneRLq1Y0/5coOyfae3\nS/YuV2hQkrboeHzMPpZ/9GTojRykr8JEX+1xzvuSE72JdDtaUkku+/z56nyFkZijcZwcBaiuiW82\nu88dyfYecc6XJid6EzVNsn/WIF0lJfjemrs0GGkcJ3keWw44f66EeJnvYY00Y1i840RvorMzfV29\nSm3uWaEbeXo4lSeFzTL5fzmqQwAAFMdJE1IsPvyJS1dHihO9RYrKnP+Bihc6fGk66Lxd3Fsmndhm\n4eFEb6JEqeuV0OAoNF7wlyaT7Sh0Tv/1eMGJ3kTyCcRvNx5UGIleFu/g+UEZi4W2if4exSMrj0sD\nZBhjTCVtE/1F/Tsq3b9OTTf3nK9POQKPR6Cy2jn9n1n8SXZg7zptE337FklK93+iQp+h5m2aqX0u\nzXTPB6vR+7FvVYfBHGxERopt45DQAAAP/0lEQVTqECKmbaLPVDyE/qMVe5XuHwCW5Rbj2KnYe6z0\n6ahPOQK5Imcgq/ccRWFZTb2kvOIT2HrwmPe6nIISnvjdZMXHTyM774jfup1Fx5Fb6N9Xfl9JOTbu\nK7UztKBuHpWuOoSIaZvo48G/f9iBlxfkwmNj37pTldV4c9EuFB8/jV+8tQL3fbgm7NsWlp3CRyv2\nYvq6fX6V+UZl2lsKVqVrXl2GCf9ejI9X7sUF//gRl72wGFXVHryxaCeueGkpzpk6HwCwNLcYq/KO\n4MdthVi79yg8HoG3Fu/yK14WSKdWTe14GA2qqPLg9YU7UVHlwdq9R7Fwe1HYt619LuoWR5uZsx+5\nhWV4d+luv+6w+0rK8Vl28JHNV7+yDNe+thwAsHbvUfy4rRAX/XMhxj2/CDkFJZi/teaLefTU+fjZ\nf5ZE8jAtk5XeNub7KK+o+ZzaNc+DNlPPeDwCT82Kr+Hd//phOwAgrW0zTBra1bs+p6AE87YU4qGL\n+9S/zffbMa5/R6wvKEGnVk2xfNdh/HJUOnq0T663raz28R8+XoEZ6/fj6dlbAADbDvofGX2zfj8q\nqz2YMKgznpy5GY9c0hdtk2uaZn71/mqsz68Z7Zcg9SCSu42q4vEIzN18CMfKK9EsqTEqqz3o3aEl\nLn9pCe45PxNZ6SnYdvAYvl63H5/ePQpNEhvj6Vmb8djEASAAf/hfDhrVGZ5a7RH4+evL0T2lOVo0\nScB9F/YCABQfr8CjUp2fXnWaenYVHcdNb63wW/faL4fjb7O2YM/hk3jqyuDTLmamNvw6WmVItzbe\n1/b95Xn4+7dbMfW7rd6yHXlTJ9a7zbLcYmw+cAx3nVdzjmZfSTkufO5HVFR78MzsrX63uf+jtd7l\nZ2ZvwY6nJ+Dnry/Hyt01R+vjB3byzvgmhMA/5m7DxEFdvGUE/jJ9I6Yt9x8IdcVLSwEAL/8ivooF\n1n5eAOCNRTuRf6Qc+UdP4qphNZ/xhduK8OXafbjl7HQ8ccUZoDrvOyEE+v/5OwA1k8JcOawrrKZN\not9eWIZ3l+apDiOg05UelJ2qxJlP/4BTUkXIF+fvCFgf59/zdvhdfnvJbr/LVw/rioXbi3DYmIU+\nKaFR0MJfhWWnkTFlFgAgJTkJR4zbPPzZegA1TUy1H9hiqb+6XJkvsbH68fvPf78dLy3IDXjd6wt3\n4XXs8l4e8bcfkNGuOfIOn8THK/MxcXBnzMo54Hebp2dtxpuLa57X1Xtqys3+sKXhZp1aY/+5sN66\nk8Y5mdreViUnKzD0ye8BAA+N832hN09S85Eb2KWVN9HXxii/9zKmzMKNI7thQOdWeHz6JvxxQj88\nM3srAOCu8zKxKu8IrjOOvGsNe3Iufnj4fLSrU2+mslrgzUW7vEkeAP7wRQ5+OSod9320Bh6PwLFT\nVX6f17pJXnbfR+H/KrVb7XMEAD9u8/9l9P7yPXjfeFzfPzQGvY2JxcukX30Pf7bOlkRP8dA7JCsr\nS2RnZ0d12wOl5Tj77/PrrQ90hGKH2qSqE34uzTP/9+crOX80fd0+/O6Tdbbv12o6vDdjeQxEtFoI\nEXIKMfW/yWMUKMkzPeg4qXV7RfO3qvolYaW2zRNDb+QAdpThdnSiX7wj/JNIzHneukW/uU5bNVWT\nnOLgNIvpXrlphOoQTNHnT9Z397Xs5Seiy4hoGxHlEtEUK/Yx9dutoTdijhXqBDQLX90T0Tro2YHf\nH+GyJNETUWMALwMYD2AAgBuJaIDZ++E+zXqrjoPzR7oYYUKXwHjTPlndpCNOY9UR/UgAuUKIXUKI\nCgCfAJhk9k7K43Ci474d1U1tphsu7WueloqajKzUqJF+v1KsYlWi7wpAHiVRYKwz1SmFk1cH8/rN\nerQbxoMkHRuWGVPAqk9SoK9av+MzIrqbiLKJKLuoSJ+Tqk0TnVfwKF6lKuqhwphurEr0BQC6SZfT\nAOyXNxBCvCGEyBJCZKWmppq6c3m+Ubu1aqZfNzbGmLNZlRFXAehNRD2IKAnADQBmWLSverq2aWbX\nrurRsb8yY8zZLEn0QogqAPcDmANgC4DPhBCbzN7Pg+N6B1zPbbuMMeZjWUYUQswWQvQRQvQUQjxt\nxT66BDly13FEpSo6TTrCmFluOLNb6I3CZEfnIUcf+gabXOTBcfWrQrLoZPKgJdPcdk6G0v0vfOQC\npfvXyV8uP8O0+/r1BT1Nu69gHJ3oL+zbAc9dO9hv3X/vHIlze7dXFFGNyZf1a/D6QV1bo0PLJkG/\nyW8fnYGv7xttQWSR+3mWeUcu0UgI83DntxcFbsardc3wNLx3+5kh7+dPE/uHtb9Arg5RhfD3l6g9\nAElv5/vSHtMnFVcO7RJwu2uGp3mXB3ZtZXlc0QjWbGuXZkmN8ew1g0NvGIaHbDgwdXQbBxHhuqxu\nyExtgWteXYZBXVvjvN7m9uCJxsgeDU819sxVgzAorTWueGkJcgr8Z815/GcDcOe5PQLe7tvfnYfu\nKc1RUl6JAyXlWF9QiqdmbsZTVw5Eq6YJ3uqE/7xuCH7/+Xrv7bY8eRk27S/FwK6tcfUry7D5wDHk\nTZ0YtALfb8f2wovzczEkrXW9Wtp269ymKfKPlCOzfTK+vn80Dh+vQErzJAx5cq53mzkPjkHfTi1R\ncrIC7y/fg7kPjUFa22b4buNBbznmX43pgX6dWmHlHy/CyGfmBdzXyIwU3HVeJv42a4vf+q1P1Tx/\np6s8OF3pQZ9OLXGqshrtk5tgz5ET3rrpz18/FFMm9MNd07K9r2vP1GR8ce852FdSHheDljY/eSlO\nnK72dl19+OK+WJxbhMe+2ojzerfHnyYOQN9OLfHctYOxr6Qc3VKaY/Weo/hm/X68tywP7ZKTcPhE\nBWbcPxqnqzze0sX/vXMkzsxIwZo9R5HePhmNqKbg4M8Gd8afLx+A33ywBtlGOWjZjSO74/PsfFR5\nBM7qkYLnrx+Kg6Wn0KFlE2w/VIbKag+SEhohrW1zbNpfioc+XY+HxvXBA2N72fq8BTK0exvv8rUj\n0vDA2F4gEBo3Jjz4yVqsyjuKj381Cm2TE9G6WSIqqwTGPLfAe5sbzuyGv15xBhLsOKcohFD+N2LE\nCBGr6mqP8Hg8Md+PGUrLK0T65Jn1/oY/OVe89mOud7u5mw6K9MkzRWl5hej/+LciffJMUXKywnu9\nx+MRj3y+Tlz9ylJRXR34scnrP8/OF73+OEscOlYu0ifPFAu2Hqp3O4/H41039Ik5In3yTJGdd0Rc\n+q+F4okZm0SV8TzGy/M5euo8kT55psgrPu63vvDYKZE+eaYY8sQc7zr5sdWqrq6/btw/fxTPzNrs\n3b662uN93EII8esPssUd764U/5y7TYx5dn6D8VVWVYs+j80Wn67cWy+OeHkOQyk4elKkT54p1u49\nGnQbj8fj996o9fXaApE+eaYor6iqdxv58d/89gqRPnmmmL/lkPB4PGLTvlKRPnmm2FlYJj78aY/o\n89hsURXkPV73PuNJsNfZ4wn82v9l+kYx8cVFpr03AGSLMHKs4+vRx7PHvtqAkT1S8M36A/jNhT0x\nvHvweiMrdh3GO0t349WbRvDQbsnFzy/EjsLjWPyHC9Etpbl3vccj8OsPV+POczND/oJi6uUWHsdf\nZ2zCm7dkoVkSDyo0S7j16DnRs7iWf+Qkvl67D/eP7aW8GYmxeBNuond0Gz3TX7eU5nggxIlWxljD\nHN3rhjHGWGic6BljTHOc6BljTHOc6BljTHOc6BljTHOc6BljTHOc6BljTHOc6BljTHNxMTKWiIoA\n7Iny5u0BFJsYjhPwY3YHfszuEMtjThdChKzkGBeJPhZElB3OEGCd8GN2B37M7mDHY+amG8YY0xwn\nesYY05wOif4N1QEowI/ZHfgxu4Plj9nxbfSMMcYapsMRPWOMsQY4OtET0WVEtI2Icoloiup4IkFE\n3YhoARFtIaJNRPQ7Y30KEX1PRDuM/22N9URELxqPNYeIhkv3daux/Q4iulVaP4KINhi3eZHiZOYO\nImpMRGuJaKZxuQcRrTDi/5SIkoz1TYzLucb1GdJ9PGqs30ZEl0rr4+49QURtiOgLItpqvN5n6/46\nE9FDxvt6IxF9TERNdXudiegdIiokoo3SOstf12D7aFA48w3G4x+AxgB2AsgEkARgPYABquOKIP7O\nAIYbyy0BbAcwAMCzAKYY66cA+D9jeQKAbwEQgFEAVhjrUwDsMv63NZbbGtetBHC2cZtvAYxX/biN\nuB4G8BGAmcblzwDcYCy/BuDXxvJvALxmLN8A4FNjeYDxejcB0MN4HzSO1/cEgGkA7jKWkwC00fl1\nBtAVwG4AzaTX9zbdXmcAYwAMB7BRWmf56xpsHw3GqvpDEMOTfDaAOdLlRwE8qjquGB7PdAAXA9gG\noLOxrjOAbcby6wBulLbfZlx/I4DXpfWvG+s6A9gqrffbTuHjTAMwD8BYADONN3ExgIS6ryuAOQDO\nNpYTjO2o7mtdu108vicAtDKSHtVZr+3rjJpEn28krwTjdb5Ux9cZQAb8E73lr2uwfTT05+Smm9o3\nU60CY53jGD9VhwFYAaCjEOIAABj/OxibBXu8Da0vCLBetRcA/AGAx7jcDkCJEKLKuCzH6X1sxvWl\nxvaRPhcqZQIoAvCu0Vz1FhElQ+PXWQixD8A/AOwFcAA1r9tq6P0617LjdQ22j6CcnOgDtUM6rgsR\nEbUA8D8ADwohjjW0aYB1Ior1yhDRzwAUCiFWy6sDbCpCXOeYx4yaI9ThAF4VQgwDcAI1P7eDcfxj\nNtqMJ6GmuaULgGQA4wNsqtPrHIrSx+jkRF8AoJt0OQ3AfkWxRIWIElGT5D8UQnxprD5ERJ2N6zsD\nKDTWB3u8Da1PC7BepdEAriCiPACfoKb55gUAbYiodqJ6OU7vYzOubw3gCCJ/LlQqAFAghFhhXP4C\nNYlf59d5HIDdQogiIUQlgC8BnAO9X+dadryuwfYRlJMT/SoAvY0z+UmoOYkzQ3FMYTPOoL8NYIsQ\n4nnpqhkAas+834qatvva9bcYZ+9HASg1frbNAXAJEbU1jqQuQU375QEAZUQ0ytjXLdJ9KSGEeFQI\nkSaEyEDN6zVfCHETgAUArjU2q/uYa5+La43thbH+BqO3Rg8AvVFz4iru3hNCiIMA8omor7HqIgCb\nofHrjJomm1FE1NyIqfYxa/s6S+x4XYPtIziVJ21MOBEyATW9VXYCeEx1PBHGfi5qforlAFhn/E1A\nTdvkPAA7jP8pxvYE4GXjsW4AkCXd1x0Aco2/26X1WQA2Grd5CXVOCCp+/BfA1+smEzUf4FwAnwNo\nYqxvalzONa7PlG7/mPG4tkHqZRKP7wkAQwFkG6/116jpXaH16wzgCQBbjbj+i5qeM1q9zgA+Rs05\niErUHIHfacfrGmwfDf3xyFjGGNOck5tuGGOMhYETPWOMaY4TPWOMaY4TPWOMaY4TPWOMaY4TPWOM\naY4TPWOMaY4TPWOMae7/A4ODDdG32pZWAAAAAElFTkSuQmCC\n",
      "text/plain": [
       "<matplotlib.figure.Figure at 0x7f582f18cfd0>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "# This allows us to plot right here in the notebook\n",
    "%matplotlib inline\n",
    "\n",
    "# Hacking up a noisy pulse train\n",
    "from matplotlib import pyplot as plt\n",
    "\n",
    "n = 100000\n",
    "noise = np.random.normal(size=n) * 3\n",
    "pulses = np.maximum(np.sin(np.arange(n) / (n / 23)) - 0.3, 0.0)\n",
    "waveform = ((pulses * 300) + noise).astype(np.int16)\n",
    "plt.plot(waveform)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "现在装饰此 `zero_suppress` 函数，使其在 CUDA 设备上作为向量化 ufunc 运行。如您遇到问题，请参阅 [此解决方案](../edit/solutions/zero_suppress_solution.py)。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 39,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "from numba import vectorize\n",
    "@vectorize(['int16(int16, int16)'], target='cuda')\n",
    "def zero_suppress(waveform_value, threshold):\n",
    "    if waveform_value < threshold:\n",
    "        result = 0\n",
    "    else:\n",
    "        result = waveform_value\n",
    "    return result"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 40,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "[<matplotlib.lines.Line2D at 0x7f5821c1e630>]"
      ]
     },
     "execution_count": 40,
     "metadata": {},
     "output_type": "execute_result"
    },
    {
     "data": {
      "image/png": "iVBORw0KGgoAAAANSUhEUgAAAXoAAAD8CAYAAAB5Pm/hAAAABHNCSVQICAgIfAhkiAAAAAlwSFlz\nAAALEgAACxIB0t1+/AAAADl0RVh0U29mdHdhcmUAbWF0cGxvdGxpYiB2ZXJzaW9uIDIuMS4yLCBo\ndHRwOi8vbWF0cGxvdGxpYi5vcmcvNQv5yAAAIABJREFUeJzt3Xl8VdW5N/DfQyaGAIEkROYAAjLJ\nFBSh1quICg44vQ6v3mK15V6LHazv9dX2tldbbdXb2ta3falaUbROVayiUlBxqgpKQOYxzIFAwhQI\nQ0jIun9k5+x1Ts68h3XW2s/388kn++yzc/Zzpid7r73Ws0gIAcYYY+ZqozoAxhhj3uJEzxhjhuNE\nzxhjhuNEzxhjhuNEzxhjhuNEzxhjhuNEzxhjhuNEzxhjhuNEzxhjhstWHQAAFBUVidLSUtVhMMaY\nVpYtW7ZfCFGcaLuMSPSlpaUoLy9XHQZjjGmFiHYksx033TDGmOE40TPGmOE40TPGmOE40TPGmOE4\n0TPGmOE40TPGmOE40TPGmOE40TPGAumjjdWoPHRcdRi+4EQvWV91BG+t2K06DMaYD7797FJc9vt/\nqg7DFxkxMjZTTPlD85u+teYYpo7ojsFndFQckb6OnmzAYws24idTh6BdbpbqcLTW1CRw79xVqDx0\nHKWFHfDzK4eifS5/dZ14rXwXAKCuvlFxJP7gI/oo/rBoM66f9YXqMLT2m4Ub8cKSHRjy8wWY+eJy\n1eFobUXlYby+rBJLth7EK0t34bkvtqsOSWvVR07iP15fpToMX3Gij+FofSPufnWF6jC0NWexXYLj\n3dVV+MXb6xRGozchRNjt+oYmRZGY4ZxfLQq7vb7qiKJI/MOJ3lJRXddq3d+/5vb6dEQmJgCY/fk2\nBZHo791VVbhu1uKwdRv3HlUUjf5K73u31br75pp/dM+J3jJ3eaXqEIzxIB+9u2bmS62bvRas3Yva\nEw0KotHb4eOnoq5fWVnrcyT+40RvmfXxlqjrq4+c9DkS/cVqQ265AMacG/nge6pD0M6oX7yvOgRl\nONEjelNDi8j2PBbf/NVVMe8L2gUwpo94OcAEnOgBvLGc2+Ld8m6cRM9Ss/NAMAbzZILXlpnddMuJ\nHsA9r61UHYIxjp4MRr9kP1z0249Vh2CMkw2n496/bPshnyJRgxN9EpqazD6tc9Onm2ri3n/iVPwv\nHLM1JvjcPf3pVp8i0d+jCzbEvX/DPrN7MgU+0e86mPj0+O1Ve3yIRH/JtHMO+fkCHyIJhofnr8cm\nwxOUW55fHH9q1ZW7DvsUiRqBT/TnP/ZRwm3um7vah0j0l+z1LNMvfPnp4XfXqw5BC6cDflYe+ESf\njBMJ2vdYs9NJJnDO8+4JegJz07+/sEx1CJ7hRM9cc6oxuaH5TZzpE2o4ndxr+VnFfo8jCY4Fa/eq\nDsEzCRM9EfUmoo+IaD0RrSWiH1rruxLR+0S02frdxVpPRPQEEVUQ0SoiGuP1k2CZYdh/LUxqu2U7\nzO7h4Ian+EKra+ob+Yw8mSP6RgD3CCGGABgPYCYRDQVwH4BFQoiBABZZtwFgCoCB1s8MALNcj1oB\nbld2z41PLVEdQsb774UbVYdgjMH/yR0AEiZ6IUSVEGK5tXwUwHoAPQFMAzDH2mwOgKut5WkAnhfN\nlgAoIKLurkfus9cNH1DBGDO3K3VKbfREVApgNIAvAZQIIaqA5n8GALpZm/UEIBc1qbTWRT7WDCIq\nJ6Lympr4fa/9tP2Ry6Ou5+H7qRvZuwBj+3ZRHYYxNj50meoQjDG8Z6eo6980dIa5pBM9EeUDmAvg\nR0KIeAWcKcq6Vv8mhRBPCSHKhBBlxcXFyYbhi1jJnqXm8RtGYu6dE1SHYYSC9jnIy+aZutzyzvfP\nx9KfXtxq/RFDq4ImleiJKAfNSf5FIcQb1up9LU0y1u9qa30lgN7Sn/cCwCOODBdtiHm/wg7Nv4s6\n+B2O1qJNbzdhQCEA4JnpZa3u45LFqXlz5kQAQHHHvFb3PWvo7F3J9LohAM8AWC+EeFy6ax6A6dby\ndABvSeu/ZfW+GQ+gtqWJJ9PUHo/9BTmzW36rdYnqZQTZ2ytb/y9v06b55G7Rjy/wOxyt/ezNNa3W\nPXrd2QCASUNKWt338cbqVutYbKN6F8S8b4ehheSSOaKfCOBfAVxERCusn6kAHgEwmYg2A5hs3QaA\n+QC2AqgA8DSA77kftjtG/iK8pvez3x4XWs5u07oFalUAJihIV7xrGG3aEC4YFN48x72YYos2s1nH\ntjkxt284za9lLMdPxS+yN7pP7KRvkmR63XwmhCAhxNlCiFHWz3whxAEhxCQhxEDr90FreyGEmCmE\nGCCEGCGEKPf+abjjwsHdQsvR/utvrua6Iun65bThYbd/+94mRZHoZ0TPzmG35//g/LDbcwxtbnDD\nwgSDoF6441yfIlErsCNjE1VRfHDasFbr+CA0eUvunxR2O79tdtjtP35U4Wc4Wnvxu+HJqHfXdmG3\nV+/mM81Y7k3QWy4/Lzvu/aYIbKI/lWCIeV52FtrmhL88fNEreWd0bht2u2uHXEWR6K9TRLNNvGYc\nFo6btZoFNtFHa4OP9OtrR4Td5tGKjOlt9m2tey0FQWATfeSkDteP7dVqm2tGt17HWttfVx92++xe\nnaNu16tLu6jrmXN8cTs5F53VutfStFE9wm6bODo2sIn+7ldXhN0eUNy6OyUATDqrW9T1zBbZ7XRY\nj+ijDm85t2/Y7eojJz2LSVfJzsD1SMTZ5qL13MUyXX+4aXTY7SXbDiiKxDuBTfQfbgj/YrTPjT7q\n8ImbR0ddz2yRNdH/8/KhUbebPiE80Rt44OTY68t2hd3OzY7+Fb1mTHhVkcam5MoaB9nAKGNjolm3\nJ97Afz0FNtFHuvmcPlHXd4i4Kt+YZJ3wIFmzO/yLEfmatWiXE/7PVLSujBF4ka9I94iL2i0iyyFw\ny01rW2vqwm7/x6WDk/q7RPPL6ogTvSXWkVOkd1dn5CBfpWa+tDyp7ZoHWdv4iL6199ftC7v9yozx\nSf3dnS8m9x4EyatLw8+O2ucm15XSxJ46nOgB3DN5UNLbRhvqz9LDE5C09s/N4TNGde/MF7DT9WTE\n5C3RatsEBSd6AAOSbLtjif37BQOS3vaX76zzMBLGwg0+o2PM+974ntlVVgOZ6PccPhF2u0/X9kn/\n7Ypdh90Oxyj3JtkOCgA1R+sTbxRgHWJ0EGDuGx2n0JkJApnoI0e49u4SP9FPHmr3vd1fd8qTmEzR\nJomBaCw5v7txVNz7+xdz+We3RF4/Mk0gE/2JiH7fndvHH1J+y7nRe+Sw1N0/5SzVIWjjGwOL4t7/\nBk/qwpIUyET/fIrV/s7p19WbQAxw8FhqZzh5SfZuCqLI8QiJeokUtOf6Qcl64MroYzuCIpDfuoYU\n+/Ul2y0riP6UYhXKG8fx2VEsn2zi0a1eSXTWHsm0SYYCmeiXbbe79cWaJJgl55nPtqW0fbuIC4y7\nDpo5o086Ivt9p4oH88V25dk9Em8kqaiuS7yRRgKZ6PdKNVaeve0chZGwyJpDQeZ0GruKGrOSkxOR\nRd6ys1JLdT8y7HMZyEQvizzCTMbeWi7GFc3gktj9lGMp50FTIRv2pj6DmTwV3qP/MG/ofrrSOSKX\nu1nzEb3m6urD55DMTfE/PdC6hgZrdn6CXiIseRdHmQQ8mouk6S8/2ljjVTjamf359pT/Zs7t5p7d\nBy7RR5aBTbbGjexnb61xKxyjlJV2SWq7yDlQWWu3f6M0qe1uGNfb20A09fJXO0PLOVnJ9ZE3uaZ/\n4BK9G7bUHFMdQkb65qDipLa7dFhyR6tBNijJZrCSTtGrWzLbreP7Jt4IQGG+ubVwApfoDR8Ap1Sy\n3VCnjeqZeKOAKzI46fitX1FyI4g7GjxReOASfbqGdudumG7pnUJtIcacunp0cgcWJpfvCFyiX7I1\nvWnC/nTLGJcj0d+xiAvbjGWiTm1TGyxlosAl+rte+jqtv0v29C9IVlXWqg6BMc9s22/OtbjAJXrm\nnk83c3e+TPPXJTtUh2CMhWv3qg7BNYFO9Gf34m5+Tsz6eIsrj3PkZEPijQy30+Go2BbPfp5aSQoW\nbt5dE0PLn1fsj7OlXgKd6GffNk51CAzAFxXpXTcxyf/7cLMrj8Ndf52NXM+Xet5ETuuos0Aneu7C\n5p6XvntuStt3kaoJ7jzIyem1ZZWh5YevGZ7S37bNCfTXuJXDJ9KfHMjUIVP8CUnTmt18IVI2YUBq\n5Q8euGpYaPm3721yOxytpVoz6M4LzvQoEj0R7G6SV45MrWqlqTjRp8m0etV+k/vS1zdyeV3ZWSmO\n2WjPc8uG2X3Yvt7x62tHpPS3PQvauR1ORuBEn4JvnGkftS7dzlUXmTfyUxyheUMZ17uR3f5ceWj5\n9OnUGmPa5pj5TzNQif6ow94dL9xhV7d7dAGXhHUip02gPnqeSnX2pCDp1M7csgapCNS3bbXDdnXT\nZ4r3E8/sxfzA39lmgUr0DSmexjHv8BeQMf8kTPRENJuIqolojbTuASLaTUQrrJ+p0n33E1EFEW0k\noku9Cjwd9XwBlTEWQMkc0T8H4LIo638nhBhl/cwHACIaCuAmAMOsv/n/RJQxVzfkyZd/f+MohZHo\njyeiZkHw1baDqkNwRcJEL4T4FECyz3YagFeEEPVCiG0AKgBkzPxcizZUh5bH9k1uNiQW3Qnp7OiC\nJCccYdF9YdBQexPI3VVveHKxwkjc46SN/i4iWmU17bRkzZ4AdknbVFrrMg7XRHfmkt99Glr+6eVD\nHD+eydO4JfLO6qrQcv80q6ROGX5GaHnzvtQnGWe2h65ObWSyDtJN9LMADAAwCkAVgN9a66NdYYv6\nDSaiGURUTkTlNTVcBVE3VVI9kWwXJmxYvvOw48fQ1dGTdl3/f7ugf1qPkSNNcn/sFF+LcuKKs80b\nTZtWohdC7BNCnBZCNAF4GnbzTCUAefRGLwB7YjzGU0KIMiFEWXExn/rrrG+h81r99Y3BTU7vSeVw\nx/btmtZjDOyWH1recSC4tYPcqISam21eZ8S0nhERdZduXgOgpUfOPAA3EVEeEfUDMBDAV85CzCyX\nDbNPkbkMQrOsNI/oz+tfGFpet+eIW+FoRy4BMaA4vX+aWVn2e/CHD9yphKmj57/YHlqWm7OCLpnu\nlS8DWAxgMBFVEtEdAB4jotVEtArAhQDuBgAhxFoAfwOwDsACADOFEEZlw5JOdsXLZTu4DIIT0yeU\nhpafk76gQZbu+AK5EutWg2ZGStVvpAJ5My/kYm8tEo4PFkLcHGX1M3G2fxjAw06CymTyRdzy7Ycw\n8czUqjYym1yq+BQXNnPk+jG9cO/rq1SHkVGG9+SJhVqY1xjlsRvG2Zcg3lq5W2Ek+hvZuyC0XH20\nXmEk+mvjwgVxZq7AJPqq2hOuPI48o/zWgM7m49a1CVMrBTKWaQKT6KuP8BGjW07xqFjGtBKYRC+X\nFT63X3pd2BhjwWPCYL7AJPr9dfYRPZc/cEb+3P/1jtTmimVMNyt26T+YLzCJftO+utDytWN6KYxE\nfw1S003XDrkKI2GyXl3safBON+l/FJopnv18u+oQHAtMopelOyiFNas8ZF/Yzs3m3h6Zom+h3fX3\n0PFTCiPR39KfXhxanrcy6uB+rQQy0fOkF87MkQY3uVH+gLmjoJ19drVZOoNlqSvumJd4I40EMtEz\nZ/7+tT1+wOl1KjcKounsk012Qb+5d57n6LEmDelmP9bySkePxczCiZ454jRPD+luzx37xZbg1WVf\nsMYuUdy5nbPrHXIFy20BLIPw5dYDoeWJZxbG2TJ4ONGnQb4AGfSLXtlZzj5C8vWS784pdxqOdk42\n2Be2cx2+ljlSYbMg1mFattN+zs/fzr3BZJzo0/Cra0aElk3oY6vS1BF2IdQg1lGvq7dr0Rd1dHZE\nH/BjDuw/al+ATreiqqk40aehtMju3dAY9G+XQ0EvCrdhr12euX1uwhqDcZWVBnt8yOzPt6kOIWMF\nItG7fdTdJFUAeG/dPlcfO2iCfuS166A7NZgAoFvHtq49FjNLIBK9PO2dG5qkfxx7XSqWFlR5Bs7m\nw1imCcS3rI3L/eblqcZeK+dubE7wmAbGvBeIRL/3iLtH9INKOoaWN1fzwBTGWGYLRKJfu6dWdQiM\nMc2M7iNPjOPuwaLfApHoCXbzwHVc0Mw1ndvlJN6IMU1dNNgeaaz7fBaBSPQvLNkRWr5qVA+Fkehv\n876joWV5gA5zZlBJvuoQWISxUnfVbM0/64FI9FukdvQLBhUrjER/h443hJb/zyWDFUZilu+c39+V\nx+nU1u6Lr3tzQyq8mFxebgnQvSZTIBI9T33nHnmU/v8q6x17Q5bQnsN211y3JsP5y/RxoeWG08EZ\nzCeXIrlypDtn7UN7dJJucaJnASJ3h3RrsJNc4+VEgMogyJOsF7o0gcs50jSZmh+EpkQe2zK6d0Gc\nLZMnX4P6q9T8qyNO9Cwlsz7e4vpj3nPJoNByU4BqBx085u3kIGt2H0m8kSHkf5peeE6ag0FHnOjT\n9L/P7RNaPn6qMc6WZnnfg5IPcmGzICX6+av3hpbJg6aBf3shONVAt9TYZZk7tnVWM8hEnOjTdP1Y\nu5vmJxtr4mzJEmmfmxVaPlAXnCnwvC7CFaR6e3I9q0uHn6EwkszEiT5NWVJb9dLtwav97abCfHva\nth0HjyuMRKEAtad74e1V9ryundry+I5InOjTJNfPCVJzg9eWbT+oOgSmob8u2ak6hIzGiT5NPQrs\nkrCNTdx90y1LtgYz0bfN4a8i8w5/utIkNzfw0YR7vgroEX1edlbijRhLEyd6xhgzXKAS/X9dOVR1\nCIwxjUweWqI6BFcEKtEP79lZdQjGOLMbF+Fi5pN78Hg9KMtLxif6L7bsDy1z5xj3fP+iM1WHwJjn\nbh1vD4zctv9YnC0zm/GJ/lfz14eWuRukM/IRTc+CdgojMctlw9wd4NOjM08S7hZ5lK3bU5L6yfhE\nX3fSLk/Aed6ZIyftEsVuVVtkQJ7LXStvGd83tHw6AMNjvShR3EJ++XQuEpfwE0ZEs4momojWSOu6\nEtH7RLTZ+t3FWk9E9AQRVRDRKiIa42XwyZDrp8t931nq5H+Ubk/qLZdBCJoh3Tsl3igFN42zy0fv\nr9N7ZqRkeFmGXG4F+HrnYc/247VkDiWeA3BZxLr7ACwSQgwEsMi6DQBTAAy0fmYAmOVOmOmrPWEn\n+r6FHRRGoj8vS7U+cNUwzx47E8m1WeTSwm6Qx3h4XSEzE8iHHCNdKlHcoqSjfXC4p/ZEnC0zW8JE\nL4T4FEDkKJZpAOZYy3MAXC2tf140WwKggIi6gxlhq4cXo+Sa9EHw2rLK0HI/Dw9AXpf2EwSPXXe2\nq4/XRZon4I3lu119bD+l++0qEUJUAYD1u2UW3Z4AdknbVVrrjDS8p7un3Jmu0sOCY210bgBNw/oq\nu1Z8F5cmHYmm5qj5TTcyL5tnd2pccM/tw6ho39aoV4OIaAYRlRNReU2NnmV+fzJliOoQfLWystaz\nx754SLfQspcX1zLFs59v92U/pwPWA6FDLteijybdRL+vpUnG+l1tra8EIE8k2gvAHkQhhHhKCFEm\nhCgrLtZzwu6C9t4diQVNe+kLumwHl312yz836XkQlYoNe+2zo6CdGSYr3UQ/D8B0a3k6gLek9d+y\net+MB1Db0sRjog55dk+RIByF+mXNbu/OHILmyEnzZz+7btZi1SFkvGS6V74MYDGAwURUSUR3AHgE\nwGQi2gxgsnUbAOYD2AqgAsDTAL7nSdQZQu7Fc6ze/C+UXx6WBrkxxpxL2KAlhLg5xl2TomwrAMx0\nGpSOgtUSyhjTSWD6tE0YUOjp49/ztxWePn4mCVpXSMZ0F5hvrDyZtxc+CtAE4Z3acc8GxnRidKKX\n280nDCjydF9BmgrO7fIHjGUyE8pzGJ2dGqWKRF7nppMNwel1k8WJ3jVcBTTzPXT1cNUhOGZ2ovew\n2FGQnT/Q27Mj08l1bkyZwchkk4bo/x4Znej/79zVoeVuHfPibMkSaZLOjrrm80Axt8j1zt308yt4\n2ky35GXbaVLXg0ejE/0H6/eFlrld2ZkDUhXE2yf2UxiJWUZ4NL2lV/9AgkiecOSwVA1XJ0Yneuae\nv3y2NbRc0smbwlHXjja2/l2Y7Qfs4lheNd0EYcIRv8hVFdbuORJ7wwzGid6h8/p72z8/UyxaX514\nI4e+LZ0pnDil70TMiayTkoVXZ5rt84J3RO9Vz7csKdMv3RZZsV0PnOgdmjrC3fk+M1VFdZ3n++gu\nlZiV56c1jR9zF4/q5e4EHDr4xVXe9I6R/xk//c+tcbbMXJzoHfqXwd0Sb8SSUiTNjKTzRMyJ+FEA\nr09h+9Dydg8njMkkQ3t4Pz9EvabFCznRO8RlUb2xslLf+TkT+eNHFb7u75F/bPB1f37avO9oaLkH\nj0mIiRO9Q3KeFwGb5MFLfidDP23z+Qh7wdq9vu7PT0el0e98yBUbJ3qH5MmDdx3Ud/LgTFNvcBs9\nc8/8VfZ0Fwa39jnGid4huelm1W5zmxv8xudGLBny54RnfIuNE72LHpi3VnUIxuBWMJaMZz7bpjoE\nLQQi0edm+/M099edSrwRS8pqnk6QMdcEItEPKslXHQJjjCkTiET/w0mDVIfAGNNY/+IOiTfKYIFI\n9FwKljHmxKSz7IGROnajNjbR12paZS7T3TahVHUIjPluVO8uoWUd68UZm+iX7zykOgQjFbTPUR2C\n1uqkAT5MTzpWBjU20VcfOak6BCN5XYOmczuz/5HIBc1Mf64mkStY+lGUzm3GJnp5dinmjDyrTr7H\n5W8HaH7RKxW/vnaEp4//6HXePn6QXDzEbqM/eEy/btTGJnrmnheW7Agt/+t5fT3dl4ZnxSmRz4cK\nO3g7krNnQfvEGxlCTsReyM6yU+Xfv97t6b68wImeJfTg2+tCyzlZ3n5kfnjxwNByRfXROFuyRORp\nCo8Zfm3gx5MH+7av/1640bd9uYUTvQvuvpj76btFnoh5XZV5iX6HNI3gGZ29mZKxRU62ff6weMsB\nT/elwiqplPWZ3XhQZDyc6F3QqwvXwXZLaaHZbfS7DtqJvq/Hz7V9rn095d65qzzdlwpbauxZzwSX\nwYuLE70LuMuhe+TJI3QcmJLIqdNqZijS8QJiIg2N9ucji2sUx8WJ3gXyZ6zJ9KuJPjIwz+M37+nX\nvpup1lXZk6xne3ztSHf86rjgnH6FoeV5K/cojMQsn1fsVx2C63hyGvfsreWxMskyPtH7UaJY7lv+\n50+2eL6/oNBxBCLzj8lTJLrN+ET/zYFFvu7P65GjQXLhWd72jWYsKIxP9ORz4uU8757FW83rEsiY\nCsYn+jaceLX10pc7VYfAmBGMT/QEfzP92j1HEm/EGNPaiVOnVYeQEkcVqohoO4CjAE4DaBRClBFR\nVwCvAigFsB3ADUIIZTWDx/XrqmrXjDFDNTQ1oR2yVIeRNDeO6C8UQowSQpRZt+8DsEgIMRDAIuu2\nr+SBNueUcqJ3C7+WLMgGl3QMLevW6cKLpptpAOZYy3MAXO3BPuL6eFNNaFmz9yOjjenbJfFGjBlK\nrk2k27U/p4leAHiPiJYR0QxrXYkQogoArN++95Hbc9gelMKJ3hn57IhfS/3cNK636hCM8f2Lzgwt\ny8XpdOA00U8UQowBMAXATCL6ZrJ/SEQziKiciMpramoS/0EK/vhhRWh5yBmdXH3soDkg1UiRP+he\nyg3AcPYxfQp82c8PJg1MvJHm2ub483kpLbKL0Ok2mM/RKySE2GP9rgbwdwDnANhHRN0BwPpdHeNv\nnxJClAkhyoqLi52E0UqVNDS6jW7nWBlmx4FjoWW5GqKXTD1zkJODXwm4UwCmK5RLkHhJt3Z5WdqJ\nnog6EFHHlmUAlwBYA2AegOnWZtMBvOU0SKbOdbMW+77P4o55vu/TD3JZ3aE9/DnTzPOhBIhqfqVf\nuZzKy1/pNcbDyaegBMBnRLQSwFcA3hVCLADwCIDJRLQZwGTrNmNJu2CQu2d4mWLtntrQsl/jO7ye\nESwT9CvyZw4DuabVi5oN5kv7XFwIsRXAyCjrDwCY5CQoHRXl52F/Xb3qMFgGazyt9sK2EML3kiB+\nuGz4GapDyHjm/7v3yeShJapDMIZel7mSJ7fRy0eHfnnq062+79MPfQuDMwl6ujjRu+Qun3qkBMF1\nY3qGlk2aZergcbsHU9sc/0dVfr3zcOKNNNS9M0/lmQgnepcU5eeGlk1KTiqM7WuPwF1VWRtnS708\ntkDt7FLb9h9LvBEzEid6l8hzVlYe4lmE3HKsvlF1CMbYuO+o6hBcU3uiQXUIWuFE75Isqb/++iqu\nYOmW99fvUx0Cy0APvbNOdQha4UTvErk3A58iu2f5DmWFT1kG021kqmpGJ/pfXj1cyX4/4KNQ16w0\nqI2eueeNr3erDkErRif6MkXVFk81NinZL2PMW7qONNYz6iSpelPqOdEzZqTOmtYOMi7Ry2132W3U\nPL1DUn9pxpg55CsDW6XaRZnO6ESvarT3viPmlUI464yOiTdizHBXnt0jtKxTF0/jEv06qWtjt05m\nVkFU4cGrhqkOgTHlRvTSc34L4xJ9fYM9O3tetj6T92a6kk5tE2/EktKzwN8h+90783vnlpMNel5/\nMy7Rc/da99Qet09Ns7PMq3qoyq+uHeHr/roZWt9fBXkeI51SjXGJvqLanGHeqn2xZX9oWUURLlMV\n+NxzY9atY33dn59+dsVQX/d31Ui54J6vu3bEuET/s7fWqg7BGHVSnZmifH+PCn9gcDVQv/NDD5+b\nivzkdzOYPMuUTgeVxiV65p4mhYcsJYa1K8vF2bq2z42zpbdOStewTOD3lNByTasVu/Qp+8yJ3kV/\nvnWM6hBcdVrhdad+hf5MD+eXpdsPhpb7KJwoQ6fmhmRk+Z3pJTrV2+FEz2J66F11FQL7F+eHlk2o\n77/7cGaUrj5wzKwxHoNK1I3vUHkglCpO9C4yIB+FOX5K3Wm+fKT29qoqZXG4JVM+G4+/t0l1CK7q\n1UXd9Ye6eh4wxZgjcqJfu1v/CpYZkudRmSFnFk4cPWknWJWTnS9cq0+VWk70LpowoCi03KDTeV0G\nkrsgmlAk7pON1apDAAB8te288W5nAAAH/UlEQVRg4o0y3K1/+VJ1CNrhRO+ivBz75Vy3h2eZcqKN\ndETf2KR/ov9gfWYkehPk8ZiOlHGid1Eb6TSyjcJTStM0NGZKwwfLBPl52apD0A4nehfJgyk27OUj\nerfMX63/xVjmng838NlRqoxN9Fec3V3p/pfv1GcwRaY7Kg02YoylzthEf/Wonok38tBpA9qVWxR2\nUDeSk7kjh4vSBZqxiV7R5FIhjRqNmkvkurG9VIfAHGrLJbtdc7ni1oJ0GJvoVfavBYAmgxK9CSNT\ng47fQffoeG5kbqJXvH+TjuiZ/viftXtUH0Smw9xEr/jN0KngUSKcIxiz6ZfmDU70CovaATDriN6g\np8KYY6pzSzqMTfSk+P+uUUf03MLLWIjq1oJ0GJvoVTMq0ZvzVBhzTL80b3CiV/1P16REzxiz8RE9\nCzGhEBczBx92uEfDPO9doieiy4hoIxFVENF9Xu0nFtXNDSYd0XPXPMZsGuZ5bxI9EWUB+BOAKQCG\nAriZiIZ6sa9YVF9ANKnXjTnPhDHn+Ijedg6ACiHEViHEKQCvAJjm0b6iUn0QatbIWNURMJY5VPfo\nS4dXhZ17Atgl3a4EcK7bO/lkUw0eeif6BNaq68GvrKzF5Mc/URqDW7IzoCCWKa+lKh3yskNzAPNr\n6UxbaYIhN17LG8f1xnfO7+/4ceLxKtFHywxhx4VENAPADADo06dPWjvJz8vGwJL8sHWbq+sAABMG\nFKb1mE5ddFY3fLihGpPO6hY245SOiIBN++rwk6lDlOx/9m1luP25cpw/sAgd2+o92USXDrn4attB\nvDlzopL9z//B+Rj38AdoQ2j1ndFNy2t5+8R+Svb/k8uHYM7iHQDceS2L8vMcP0Yi5MWFNiI6D8AD\nQohLrdv3A4AQ4tfRti8rKxPl5eWux8EYYyYjomVCiLJE23l1yLkUwEAi6kdEuQBuAjDPo30xxhiL\nw5PzYSFEIxHdBWAhgCwAs4UQa73YF2OMsfg8a/gUQswHMN+rx2eMMZYcva8WMsYYS4gTPWOMGY4T\nPWOMGY4TPWOMGY4TPWOMGc6TAVMpB0FUA2BHmn9eBGC/i+HogJ9zMPBzDgYnz7mvEKI40UYZkeid\nIKLyZEaGmYSfczDwcw4GP54zN90wxpjhONEzxpjhTEj0T6kOQAF+zsHAzzkYPH/O2rfRM8YYi8+E\nI3rGGGNxaJ3oVU9A7gQR9Saij4hoPRGtJaIfWuu7EtH7RLTZ+t3FWk9E9IT1XFcR0RjpsaZb228m\nounS+rFEtNr6myeIMmO2SyLKIqKviegd63Y/IvrSiv9Vq7Q1iCjPul1h3V8qPcb91vqNRHSptD7j\nPhNEVEBErxPRBuv9Ps/095mI7rY+12uI6GUiamva+0xEs4momojWSOs8f19j7SMuIYSWP2guf7wF\nQH8AuQBWAhiqOq4U4u8OYIy13BHAJjRPpP4YgPus9fcBeNRangrgH2ievWs8gC+t9V0BbLV+d7GW\nu1j3fQXgPOtv/gFgiurnbcX1YwAvAXjHuv03ADdZy38GcKe1/D0Af7aWbwLwqrU81Hq/8wD0sz4H\nWZn6mQAwB8B3rOVcAAUmv89onkp0G4B20vt7m2nvM4BvAhgDYI20zvP3NdY+4saq+kvg4EU+D8BC\n6fb9AO5XHZeD5/MWgMkANgLobq3rDmCjtfwkgJul7Tda998M4Elp/ZPWuu4ANkjrw7ZT+Dx7AVgE\n4CIA71gf4v0AsiPfVzTPZ3CetZxtbUeR73XLdpn4mQDQyUp6FLHe2PcZ9pzRXa337R0Al5r4PgMo\nRXii9/x9jbWPeD86N91Em4C8p6JYHLFOVUcD+BJAiRCiCgCs392szWI933jrK6OsV+33AO4F0GTd\nLgRwWAjRaN2W4ww9N+v+Wmv7VF8LlfoDqAHwrNVc9Rci6gCD32chxG4AvwGwE0AVmt+3ZTD7fW7h\nx/saax8x6ZzoE05ArgMiygcwF8CPhBBH4m0aZZ1IY70yRHQFgGohxDJ5dZRNRYL7tHnOaD5CHQNg\nlhBiNIBjaD7djkX752y1GU9Dc3NLDwAdAEyJsqlJ73MiSp+jzom+EkBv6XYvAHsUxZIWIspBc5J/\nUQjxhrV6HxF1t+7vDqDaWh/r+cZb3yvKepUmAriKiLYDeAXNzTe/B1BARC2znclxhp6bdX9nAAeR\n+muhUiWASiHEl9bt19Gc+E1+ny8GsE0IUSOEaADwBoAJMPt9buHH+xprHzHpnOi1noDcuoL+DID1\nQojHpbvmAWi58j4dzW33Leu/ZV29Hw+g1jptWwjgEiLqYh1JXYLm9ssqAEeJaLy1r29Jj6WEEOJ+\nIUQvIUQpmt+vD4UQtwD4CMD11maRz7nltbje2l5Y62+yemv0AzAQzReuMu4zIYTYC2AXEQ22Vk0C\nsA4Gv89obrIZT0TtrZhanrOx77PEj/c11j5iU3nRxoULIVPR3FtlC4Cfqo4nxdi/geZTsVUAVlg/\nU9HcNrkIwGbrd1drewLwJ+u5rgZQJj3W7QAqrJ9vS+vLAKyx/uaPiLggqPj5/wvsXjf90fwFrgDw\nGoA8a31b63aFdX9/6e9/aj2vjZB6mWTiZwLAKADl1nv9Jpp7Vxj9PgN4EMAGK64X0Nxzxqj3GcDL\naL4G0YDmI/A7/HhfY+0j3g+PjGWMMcPp3HTDGGMsCZzoGWPMcJzoGWPMcJzoGWPMcJzoGWPMcJzo\nGWPMcJzoGWPMcJzoGWPMcP8De27I6dEy0koAAAAASUVORK5CYII=\n",
      "text/plain": [
       "<matplotlib.figure.Figure at 0x7f582c03df28>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "# This will throw an error until you successfully vectorize the `zero_suppress` function above.\n",
    "# The noise on the baseline should disappear when zero_suppress is implemented\n",
    "plt.plot(zero_suppress(waveform, 15))"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 管理 GPU 显存\n",
    "\n",
    "到目前为止，我们一直将 CPU 上的 NumPy 数组用作 GPU 函数的输入和输出。为方便起见，Numba 已经自动将这些数据传输至 GPU，以便由 GPU 处理。通过这种隐式数据传输，Numba采取保守的行动，将在处理后自动将数据传输回CPU。可以想象，这是一个非常耗时的操作。\n",
    "\n",
    "[CUDA 最佳实践指南](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html) 中指出：\n",
    "\n",
    "> **高优先级**：尽量减少 CPU(Host) 和 GPU(Device) 间的数据传输，即便与 CPU 相比，在设备上运行某些核函数并未实现性能提升时，也需遵照此规则。\n",
    "\n",
    "秉持这一点，我们应思考如何避免将数据自动传输回主机，以降低开销。我们同时要确保仅在准备就绪时，将数据及时复制回主机。\n",
    "\n",
    "为此，我们需创建 **CUDA 设备数组**，并将其传递至 GPU 函数。设备数组经过处理后不会自动传输回主机，因此您可根据需要在设备上重复使用这些数组，并只在操作至最后一步时，再将全部或部分数组发送回主机。\n",
    "\n",
    "为进行演示，我们再次创建加法 ufunc 示例："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 41,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "@vectorize(['float32(float32, float32)'], target='cuda')\n",
    "def add_ufunc(x, y):\n",
    "    return x + y"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 42,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "n = 100000\n",
    "x = np.arange(n).astype(np.float32)\n",
    "y = 2 * x"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 43,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "1.13 ms ± 1.33 µs per loop (mean ± std. dev. of 7 runs, 1000 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit add_ufunc(x, y)  # Baseline performance with host arrays"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "在 `numba.cuda` 模块中，有类函数可以将主机数据复制到 GPU，然后返回 CUDA 设备数组。注意，下面我们会尝试打印设备数组的内容，但我们只能获取数组信息，而无法获取其实际所含内容。这是因为数据目前位于设备上，我们需将其传输回主机才能打印相应值（稍后我们会对此进行讲解）："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 44,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "<numba.cuda.cudadrv.devicearray.DeviceNDArray object at 0x7f5831848c18>\n",
      "(100000,)\n",
      "float32\n"
     ]
    }
   ],
   "source": [
    "from numba import cuda\n",
    "\n",
    "x_device = cuda.to_device(x)\n",
    "y_device = cuda.to_device(y)\n",
    "\n",
    "print(x_device)\n",
    "print(x_device.shape)\n",
    "print(x_device.dtype)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "与 NumPy 数组类似，设备数组也可传递至 CUDA 函数，但在复制时不会产生任何额外开销："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 45,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "563 µs ± 966 ns per loop (mean ± std. dev. of 7 runs, 1000 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit add_ufunc(x_device, y_device)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "由于 `x_device` 和 `y_device` 已存在于设备上，因此该基准测试会更快完成。\n",
    "\n",
    "不过，我们仍得为 ufunc 的输出分配设备数组，然后将该设备数组复制回主机，即使我们并未在以上单元中将该数组分配给一个变量。如要避免这一点，我们可以使用 [`numba.cuda.device_array()`](https://numba.pydata.org/numba-doc/dev/cuda-reference/memory.html#numba.cuda.device_array) 函数创建输出数组："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 46,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "out_device = cuda.device_array(shape=(n,), dtype=np.float32)  # does not initialize the contents, like np.empty()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "然后，我们可以在 ufunc 中使用特殊的 `out` 关键词参数，以指定输出缓存："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 47,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "455 µs ± 1.1 µs per loop (mean ± std. dev. of 7 runs, 1000 loops each)\n"
     ]
    }
   ],
   "source": [
    "%timeit add_ufunc(x_device, y_device, out=out_device)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "此次调用 `add_ufunc` 并不需要在主机和设备之间进行任何数据传输，因而运行速度最快。如要将设备数组传输回主机内存，我们可以使用 `copy_to_host()` 方法："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 48,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "[ 0.  3.  6.  9. 12. 15. 18. 21. 24. 27.]\n"
     ]
    }
   ],
   "source": [
    "out_host = out_device.copy_to_host()\n",
    "print(out_host[:10])"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "您可能会认为，此处比较的对象类型并不相同，因为在使用设备数组时，我们并未对 `to_device` 调用执行基准测试；但在使用主机数组 `a` 和 `b` 时，隐式数据传输会被纳入基准测试。事实的确如此。当然，正如之前所述，`add_func` 函数并非特别适用于 GPU，以上示例只是为说明如何消除传输需求。\n",
    "\n",
    "如要确定是否应向 GPU 传输数据，请务必对您的数据传输执行基准测试。\n",
    "\n",
    "此外，Numba 还提供其他方法来管理设备内存和数据传输，详情请参阅 [此文档](https://numba.pydata.org/numba-doc/dev/cuda/memory.html)。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 练习：优化内存传输\n",
    "\n",
    "给定以下 ufunc："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 54,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "import math\n",
    "\n",
    "@vectorize(['float32(float32, float32, float32)'], target='cuda')\n",
    "def make_pulses(i, period, amplitude):\n",
    "    return max(math.sin(i / period) - 0.3, 0.0) * amplitude\n",
    "\n",
    "n = 100000\n",
    "noise = (np.random.normal(size=n) * 3).astype(np.float32)\n",
    "t = np.arange(n, dtype=np.float32)\n",
    "period = n / 23"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "如下方单元所示，在调用 `make_pulses` 和 `add_ufunc` 期间，将数据传回主机并再次传回设备是一场无谓的数据往返之旅。\n",
    "\n",
    "更新下方单元以使用设备内存分配功能，确保在调用 `make_pulses` 前仅会在设备中复制一次数据，且在调用 `add_ufunc` 后亦仅需在主机中复制一次。如您遇到问题，请参阅 [此解决方案](../edit/solutions/make_pulses_solution.py)。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 56,
   "metadata": {},
   "outputs": [],
   "source": [
    "t_device = cuda.to_device(t)\n",
    "noise_device = cuda.to_device(noise)\n",
    "pulses_device = cuda.device_array(shape=(n,), dtype=np.float32)\n",
    "\n",
    "make_pulses(t_device, period, 100.0,out=pulses_device)\n",
    "waveform = add_ufunc(pulses_device, noise_device)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 57,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "[<matplotlib.lines.Line2D at 0x7f5821ba8eb8>]"
      ]
     },
     "execution_count": 57,
     "metadata": {},
     "output_type": "execute_result"
    },
    {
     "data": {
      "image/png": "iVBORw0KGgoAAAANSUhEUgAAAXQAAAD8CAYAAABn919SAAAABHNCSVQICAgIfAhkiAAAAAlwSFlz\nAAALEgAACxIB0t1+/AAAADl0RVh0U29mdHdhcmUAbWF0cGxvdGxpYiB2ZXJzaW9uIDIuMS4yLCBo\ndHRwOi8vbWF0cGxvdGxpYi5vcmcvNQv5yAAAIABJREFUeJztnXd4FWX2x78nnRAgEEINEJpgaIIR\nKRaKoGB3/a2oKPa1rKhYQHGtiFjWdd3VVQRddnURRRSUJl1QCYbeCZ1QQkKA0ELa+/vjzk1umdun\n3HnnfJ4nT+68970zZ+7M/c5bznsOCSHAMAzDWJ8Ysw1gGIZhtIEFnWEYRhJY0BmGYSSBBZ1hGEYS\nWNAZhmEkgQWdYRhGEljQGYZhJIEFnWEYRhJY0BmGYSQhzsiDNWzYUGRmZhp5SIZhGMuzevXqIiFE\neqB6hgp6ZmYmcnNzjTwkwzCM5SGifcHU4yEXhmEYSWBBZxiGkQQWdIZhGElgQWcYhpEEFnSGYRhJ\nYEFnGIaRhKAEnYieIqLNRLSJiKYSURIRtSaiHCLKI6JpRJSgt7EMwzCMbwIKOhE1BzASQLYQojOA\nWADDALwF4G9CiPYAjgO4X09Dteb0+QrMXHfQbDMYhtEBIQRmrT+E0vJKs00xlGCHXOIA1CKiOADJ\nAA4DGABguvL+FAA3aW+efoz9biOe+GodNuSfMNsUhmE0JmdPMUZOXYs3Zm812xRDCSjoQoiDAN4F\nsB8OIT8JYDWAE0KICqVaPoDmehmpNWfOV2BP0RkAwAvfbURlFSfKZqKHDfknsFe5P5nwOFXqkKbD\nJ8+ZbImxBDPkUh/AjQBaA2gGoDaAISpVVVWRiB4iolwiyi0sLIzEVk04fqYMnV6ejw35JwEAmw6W\nYMXOIpOtsjbFZ8qw/9hZs82Qhhv++Qv6vbvUbDMsDSn/qwRw6IR9RD2YIZerAOwRQhQKIcoBzADQ\nB0CqMgQDABkADql9WAgxUQiRLYTITk8PGFtGd8apdMFIpR4TPL3fXIQr3llithkMUw0pP+rF246i\nz4TF2Kg04GQnGEHfD6AXESUTEQEYCGALgCUAblXqjAAwUx8TtUMIgW/X5HuVEyt6RJyvqDLbBIZx\nw/M3vfeYPYawghlDz4Fj8nMNgI3KZyYCGA1gFBHtBJAGYLKOdmrCvE1HVMuJ2+hhc7asInAlJij+\nvjAPmw7aoyWpN+Sh6HZptAXl5SKEeFkI0VEI0VkIcZcQ4rwQYrcQoqcQop0Q4v+EEOf1NjZSjp5S\nN3H45ByDLZGHrJfmm22CFOTsPoa/LdyB6/6xorosc8xsnCuzl9udVpScK3fb3nfsLK58ZwkKfWiA\nLNhqpahdntKM9bht4krV8s2HuMUeDk98tc5t+/Nf9mLfsbOYt+mwSRYZg60E/Xy577FeIdh1MVSe\n+Wa92SZIz60f/2a2CVIQY5PGnK0E/V/Ldvl879+/7jXOEEmYvtp9gjlzzGw8MOV35B9nF0bGHA4U\nn8Xo6Ru8yn0Nt8qGrQS9+EyZz/d8TZgy6hSUlKqWL9x6FP3eWWqsMQyjMPKrtZiWe8BsM0zDVoLu\nj5w9xTjqQ6QYb4ZP8j2RXMErbxmTqLL5vWcbQa+oDOwrXVBij25ZpFRUViHv6GmzzZAGuwWQ0pP1\nARYQnZP8u44LXEUOHvrvarNNkIZgurSVVQKxdpmJioC+ExbjoI2WppvN+DnbsPfYWYy/uYvZpuiC\nbVroi7cdDVhnIy/qCIoTZ8sD1mn7whxM4YnmgAQj5t+u9l7dzITP/3L2m22CbthG0IMhd1+x2SZY\nAs9FG774YFGezpbYg6fZPZQJEhZ0F75bywkvgmFukB5BvJCLMZLzFXKPjwcDC7oLvLYoOPYXs5+5\n0by/cIfZJkQ96w8EP2Ta7dWfdLTEPGwh6KHE6uYVo1rCTXR/hOLd8v5CHr7yhxACf/wk+FW1J4Mc\nNrQathD0Oyapx8lQ43gQE352ZvG2gqDrVvHD0S8vfr/JbBOkYcvhErNNiApsIeihRFjr8foC/LSZ\nV4364r5/5wZd19/KXAbsy68hj3yxxmwTogJbCHqoCRiemrYucCUbwgtgtGX9AU5QrhU8r+NAekHn\nZLvaMehvy8w2gWEYP0gv6OEk2/XMdsI4OFAc+orGYEIuMMHxzvxtZpvARDnSC3o4sJxrx+NT15pt\ngjR8uGQXyvkB6YWsHivhwILO6Eqwi5CY4LBLXO9QkNWnPBxY0NXgJjoTpfyw/pDZJjBRDAs6w5hA\nuMMEi7YGvw6AsR8s6CqcKq3gZBeMrnzsJx0iYwy+sm5ZGRZ0H9w1eZXZJkQVe9j9UzN2F57Gv5ay\noJvN9f9YYbYJmsOC7gNOOuBO/zDcPxl1Xv1hS9if3VPEC2i0QkbvZKkFnVc2RgenStmtzJVIXA+L\nTrOXi1bIGGpIakF/7EuO7xANyDhWGQm/7jpmtgkMHHNlsiG1oC8KIu2cL06fl+9im8WPGw6bbQLD\neCFjwmipBZ2JDoLJQcow4fDKrM1mmxBVsKAzuvNvThbN6ESk99YxyeYkpBV0jrLIMEwgZIs1JK2g\nHz7JE3GMnKzeV2y2CdIg2wS1tIK+q5CzwTDRR2VV5L5yz07foIEljIxIK+icr1E78o/zYhat+N+q\n/RHvI0bGFTFhkFdwymwToo6gBJ2IUoloOhFtI6KtRNSbiBoQ0QIiylP+19fbWKORbcIkXC57a0nE\n+/hubb4Gllif00H4Pm997Rq/7+/kXKQAgLNl2rgdloWYojKaCbaF/ncA84QQHQF0A7AVwBgAi4QQ\n7QEsUrYtw543hwasw7Gng2PvhGsD1nlq2nqpfjjhEhfjv3W97Nl+qJUQa5A11iaYwauWDZID1tkh\nUUs/oKATUV0AVwCYDABCiDIhxAkANwKYolSbAuAmvYzUmgm3dAER4aXrsvzW48VFwXNDt2YB6/z1\np+0GWBLdFJ3x30hITU4AAPzvwUuNMEd6fn6uP666sLHfOjKFAAimhd4GQCGAz4loLRFNIqLaABoL\nIQ4DgPK/kY52akpWs7oAgHv7ZuLDO3rAV6Ppvs9/N9Aqa/PB7d0D1tl7jF1JP1m22+/7iXGOn2RW\n07pGmGMLJo3I9vu+CKqtbw2CEfQ4AD0A/EsI0R3AGYQwvEJEDxFRLhHlFhYWhmmmtjSumwTAkQz6\n2q5N8dw1HVXrneIWekBWjR0YdF3iVFABSYp3DLc4W+qMb6oCNK1vuihwrxGwXws9H0C+ECJH2Z4O\nh8AXEFFTAFD+qwZOEUJMFEJkCyGy09PTtbA5YpyC7uThK9uaZIn1aVQnKXAlhtGBiQF6Ox2D7OVI\npOeBBV0IcQTAASLqoBQNBLAFwCwAI5SyEQBm6mIhIw3sbeebVWMHYtHTV5pthqWYt9l3AvIZj/bB\nQ5e3qd6e+Vhfn3UDtfStRLBeLo8D+JKINgC4CMB4ABMADCKiPACDlO2oQPi5QG/e0sVAS+Tmz/3b\nmW2CpfCXNKVRnSS0TU9xK1vyTD+f9bceLtHKLCnp0bI+Ylwmx7q1SPVZVyI9D07QhRDrlGGTrkKI\nm4QQx4UQx4QQA4UQ7ZX/UbMe2d9ivNt7tjTOEAnw93B85uoObtvfPtLH777s3kIfP3trSPVbN6zt\n872bPvwlUnOkpVGdxBA/IY+iS7lS9IUZGzXb1zmNFi9YlYVbg48pf3Er/2vLik6VRWqOpSnRMHPT\neZv79PsL7bHgqdCGrjSIxhA1SCno03IPaLavp6at02xfVuToqdCCnM0ZebnP945w5iJGI276p+8e\nSr3k+JD2pUV8nWhBSkH3xdq/DAr5M/4mXhhvsprVxb/vvUT1PZn8fRlz0dKleNjElZrty2xsJej1\na7Nvb6isyCtSLb/iAt8uqP06qK8xO1Dse1LQzix/rr/ZJlgOredjNuSf0HaHJmErQWdCZ+4m9R7K\nf+7rabAl1seXe1zdpNCGCJyUV9p7HD1U2viZZF6ff9JAS/TDNoLepK7/BTCf36M+TMAwWvHLTm2T\nKQQTuVFWfDXQZzzq29NqxqN90LSeug7sliR/gnSCvvlQeE/afh2iYxWrFajN0QBDZh/HsdEUX/OY\nPVr69rRKTU5A7zZpqu+VlsvhzSadoJ/y0Wrp39G/YJPdnaRDINgl1UwNuwt9C3q4k8UyhX01Cl/z\naLIsLpJO0P+5eKdqeVazegZbYm+eHnSB2SZYBn9i8vL1vkM8l9h4yCVcnvVYDOfkq9+1c3U2E+kE\n3eeQSxCP4AC5BxiFVkEkDXh8YHvV8vUH5PAmCJkw7617+7bW1g6b44xmKSvSCbqvpBTBXMg5T6gv\nihn7nXYrT61E7l71aA7jbu4c9j7tGhPdn54HylDUM7OBajknQg+PFg1qmW2Cbkgn6OWV6i3xW3pk\nBPzsBY3qqJZ/mRN5Yl8rsvmQegCo5IS4sPc5LsR4JrKzcNSVARsbo4eox+ufMHebHiZZltggu9jT\nH/Yfc8jKSCfoanw8/OKgLnYMj7m48fKszV5lT16lPpQSLIU2zdP6+o9bVMvbNUpRLXclUIwcu/HL\nTvXFbv4CybnimQ9BJmwh6GkpvEJUK57wMTbO+GeXipdLrRDGc1+9oZOW5liaOyflqJZ/7yfmeTDI\n4LpoC0FPiLXFaRoCu3dqx+BO/pMXu+Iv1ALjoGuG75jnwfDEV2s1ssQ8bKF0XTPYZZGJPkLp+vuL\njc5ow9Lt0ZHzOBJsIejcqjSHj+7sYbYJUc1dvVpFvI8qiUK/RkJCnC2kLCD8LXhwJXdtNWNol6Zm\nmxDVtAjCnz8QZ8p4cREAxGvg0CDDo5EF3YNEftIzFkIGEdICLXrhZRJkgZJKvfI4tgVjM3jIhXFF\nKkH/YcPhiPfRMOQEs0yoBOsvzATmu7UHzTaBiSKkEvQPFuVFvI8Xr71Qtdxuy6zVWn5J8drcLtyo\n1I5Xf1BfsGQ32O3BgVSCrsbl7RuGVN/XsvadR+0l6Dd+6J2Ed9vrQ0Lez+hrvJet7y8+G5ZNVuV8\nhfUXrEQLp0rLVcv/cHHg0B52QHpB791WPaC9P6660HvBxxkNk9JagY0HtUnJ1amZd+z0T5bt0mTf\nVuHez3/3KmvmI3MO45+35qnHr3npOt9hhu2E9IIeznBtarJ3jsdRX6/XwBr7ofb17ztmrxb6r7u8\nU8/dqYEPuh0pLff2RLmvb2vN4jAt2X5Uk/2YhQ0EPXRFZ9dF7VD7nflKlmwnwrnHLsnkIF1qvOQn\nCUio7C2ydnhn6ZUrHO1gudGOeJU4Ojl71OOs24msMNL4/ff+S3WwxFoUnynTdf9Wb2tIL+j9OzYy\n2wRbU6+W9/AVg7DcMmTPthMMi7fpOyRicT2XW9AT4mLQuXnogbms/pTWg3AjVl7ICaWZKGTjK4NV\ny62+RkJqQZ/rI6VcIOrWCj8jj6zMeDT8LC++UqgxoRNMQgxZKfHhshgOdZLUe45HTpZqdgwzkEbQ\ni057Z8Jpmx7ezf/UVfbOWK+WlzWSaHZN2EXPi4Yp4a1I7tnavg/HuRsjXwnuym3ZLbzKrL7oTRpB\n/2qVdnk/k+JjMewS74ttF9724esbLiP6ZGq6P6vTp20aLmisnr82EG1U4qJvzNdmzQADCIuPoksj\n6FrHtEhJtO+wy39+2+dVFomXb7DJe+1ClwgSrtzbt7VX2crd3n7uTGAa1vFOTfn5L3uNN0RDpBF0\ntZyNkTAywmTIDOOL3m1CX73sRO3h+MacrZGYY1ueGCjf0GrQgk5EsUS0loh+VLZbE1EOEeUR0TQi\nkioTMy8ucodd5rSjXwd2pY0GZMxyFMoZPQHAtSnwFoC/CSHaAzgO4H4tDTObGE5b54YW2XVcuePT\nlZruj5EfNY/CN27ubLwhUUxQgk5EGQCuBTBJ2SYAAwBMV6pMAXCTHgaahdoKR6vHeTCLVioPA7X4\nJjJityideqIWpbMpe1C5EWwL/X0AzwFwRsZJA3BCCOH0b8sH0Fztg0T0EBHlElFuYaG1s2qrRc1j\nAlO/tlSjcSFhtzj6eqLmUtgrgvkIGQko6ER0HYCjQojVrsUqVVX9fYQQE4UQ2UKI7PR0TsDM2Is/\n/Xd14EpMUKh58/jKXxAsso2jB3M2fQHcQER7AXwFx1DL+wBSicj5bWYAOKSLhUFwtMTaq7uinVu6\nq3a+GJN40qYeWOsOnNB8n/f2zdR8n2YSUNCFEM8LITKEEJkAhgFYLIS4E8ASALcq1UYAmKmblQGo\nUOmLvTDUO1MOEx7dWqSabQLjwpDOTc02QRqGXdLSbBM0JZL+xmgAo4hoJxxj6pO1MSl01MZ62EtF\nO2rbeJFVNMK3tna0Vll5u3BLgQmWaENIgi6EWCqEuE55vVsI0VMI0U4I8X9CCO9gKgzDaE66ShwY\nq0cJjCYe+E+u2SaEjRQzArtVPAn0Whi09XCJLvuVnb7tvL0Ryiq804kxgVHzGvrH4p0mWGIuHFLC\nGykEXS3ofbYGIVvH39zFq+zoKft1RAZ38k6aHSqf3p3tVbZ63/GI9xvNGNlqnrEm37BjRQvcK/FG\nCkEnFS9KLa71HZd6T5icK6uMfMcWo66P2NGhoOZedujEuYj3G81MXXXAsGPttVnibQBoXr+W2SZE\nHVIIulrge70SET/7zXpd9hstzNE45rQ/fs6z9kKzQCzbwSuLteKUym/8tRt42b8nUgj69NXGdTdP\nqSR/kInvNQ5D7I8th+Sejzir0ptrXDe8xBZ2Z2+Rdw+E8wV7I4Wgq9EqTdtgUnZBT3dPz4nRPI5z\nwgTJN6uNG76yMtIKuq+cgYx/9PRxbp5qrzHP5XlFXmVq8z3hoDZhLzNqSVe0om6SPOsspBV0Jjx+\n31us276tnq9RC+7q3UqT/ahN2DMMCzrjRtHpMt32LVsgpHB4rH87s01gPLi2azOzTdAMy//CqrjZ\npyt/zM7QbF8PXd5Gs30xjFbcJlFCeMsL+tvzt3uVfftIb832f3tPeS52OLRK8451ES6ZKnEzGMZs\nLpIo+JzlBf2bXO/Z764ZWl4gey8vrsW5RC3DibP6DZdFG5Ek2pYZywu62oCLWvq4cOnfgZNy6Mlp\nyf36XXnzFn09U/Sc/4g2Jt59sdkmRCWWF3S9VoQ6Gdypia77tzs/75Bzteh+laX4t/fU1zPFTvlL\n9XZLLq+0ZuA4ywv6ibPeS4IZxmzKTBEEdhAIl2u7uicNufaD5SZZEhmWF3TG2sg6Q2FGEooVO70X\nMjHBcZ9HKrodBdbs7bCgM9UYEY70+m7y+Pz6w4wH1W+7vJMoM8Ehy8py6QR99DWcSzRc1uef1P0Y\nnusGSivkDEdMJjTRdxWeMfyYRjBznf4B4y5oXEf3YxiBdIL+SL+2ZptgWbYf0T/6YaWHoB+T1DPD\nCDmXyX/aHzPXHTLbBMsgnaAz4WNEQoYKD0GXdZhg3uYjuh/j+8f66n6MaKCgpNRsEywDCzpTjVoS\ngcR4bW8RTzfTRSrpA2VgwtxtZpsgDZslj5uvJSzoYSBr/BjPMdhRgy7AH7O1DX3QNp2X/zORMbyX\nMZEmrZizlAU9DI7YpAs4cmB7TVfdAsCj/TjaIBMZRnmkLLXgojcW9DDI2SPnuK8R1K+dYLYJpqBl\nwDi7k2hQGGYrxsZhQQ+Dp6bJlyi6wqJLna1ChyZ1zTZBGowKGKdVdikjYUFnAAAHT5wz2wSpibGe\nNkQt6XU40bYvLC3o3KpkrIKeybddkXXC3pWerRsYcpz9xd4B1qIdSwv6xOW7zTZBGqzYvWS8+XHj\nYbNN0J2M+smGHMeKvSpLC/rhE8Z4m3z3aB+vMiu6NDHmYVQLfeyMjYYcxw6YEb4hUiwt6MKgcKHd\nW9b3Ksvdd9yQYxvF0h3uC3xaNjCmFQQA6w6cMOxYZmGUNpyyUcIQrVk46gq37eNn2MvFUDwbyd1b\nGhfborRcrqBSC7YUuG3Xq2Vc9LkDFhyrDBWt/fmdvPUHfbMg2Yl2jdwDdE1ascckS8LH0oLumdwi\nzsBBL88gU1bHs3v50vVZJlkiHytG99dt31YcFmD0w9KCPneT+wSQkcPaeqe+MxrPVHCXZOrnSXDT\nRXLHRJ+xJt9tW0/RNWpsnrEGAQWdiFoQ0RIi2kpEm4noCaW8AREtIKI85b/3QLPOeDaSGxi4CnEL\nBwwKm8b1kty2ZXs4Tvltn9u2npJrRU+MUDjDcwIhEUwLvQLA00KICwH0AvAYEWUBGANgkRCiPYBF\nyrZhnFTJJZqWop+gfzzcPcv4uz/t0O1YsuPpIvnEV+tMskQf1ntM8urZiJa9hX6+gteahEJAQRdC\nHBZCrFFenwKwFUBzADcCmKJUmwLgJr2MVOPPU9d4lXXN0G9S9PL2DXXbt92QXIO80NPHX89GTDRg\ns1slYkIaQyeiTADdAeQAaCyEOAw4RB9AI62N88fyPO+EuMMu0TbUqyt2EyE9ubpTE7NNMBQ9PYb0\n8p6JFj5cstNsEyxF0HcDEaUA+BbAk0KIoAeQieghIsolotzCQn3DUfLkkzWwS+o0J7US9AsmZeR6\nATOwouugmQQl6EQUD4eYfymEmKEUFxBRU+X9pgBUU88IISYKIbKFENnp6ela2MwwjEKz1Fpmm2Ao\nfx92kdkmRDXBeLkQgMkAtgoh3nN5axaAEcrrEQBmam9e9CBzA90OAZ0YOWjTMMXQ4022WA8hmBZ6\nXwB3ARhAROuUv6EAJgAYRER5AAYp29Iic/Aqz8TNDBMNREO8pNd/3GK2CSERF6iCEGIFfE82D9TW\nnPDp1EzfBALxsfIKejSECa2sEoiV3amaCYnScnZZDBVppshHDmyv6/5lXmI97ff9bttDOhvvhVJR\nxT/ecGmYwgkftOKRfm3NNiEipBF0I+O4yMb2gtNu2+0b1/FRk4lGxt/c2W278NR5kyzRlh0Fp7zK\nEnTOJ2p1GZFG0PV0DbMbZtzTmw7KEUrBjJCrniJXUuq9itqKnFOJaNqhib6NjSsvMHQ5jeZII+iM\ntRn97QazTdCEV37YbPgxL2vnvorZ4o3Mas6VGR+i2qj0dnohjaDruezfF0Wn5ejartvvnqwjMd74\n22Ln0dOBK1mAmesOGX7MOElXi97779/NNsFySHMnpCQGdNjRHFkiwZWUup/Hg5e3MckS+TAjVPCi\nrapr/CxP5+b6erLJgDSCbgTXdmnqtn3HpzkmWaIvsscHMZLUZOODZ70xZ6vhxzSCikrz/dKjHf7l\nhsCgrMZu2wdPnDPJEutjZOx6M3nyKn3dae1EFKwzinpY0ENAYld0w0nU2f0sWqiTZFxuVtkZ5+Ge\nyXhjyV8Vxx6xPl8+cKlxx8rZh8wxs03xmrC6X3M00VFnl0UZsKSgm4XnkAsTPm3SjQuy9NGSXQCA\nY2eM90qSeYWx0SQnGO/4YDVY0EOAbyj7ca6sEq/+sBlny+TwaLIyHOsnMJYU9KOSLG1m1Jm/+Qi+\nWLkvcMUIKTp9HqcDuJ5+9ssefP7LXnyybHfA/R2TZF0Cow0Tf95leMYlywl6WUUVer25yK3sDRMn\nS85XRDYuu2BLAf7z296QPlN46jy2HtZmqfyhKPTU+dN/V+PF7zfpfpzscQsx4N2lqu/tLjyNisoq\nVCrzNVU+XCz2FJ3BrPWOxUSVFpvbKSktx72fr8LRklKzTYkq/ucxv1NRGV7guPFztuGd+du1MClo\nLCfoalH5OjerZ4IlDoL5EZdXViFzzGx8sCjP670H/5OLl2aGtlx84F+XYsjfl4f0mdPnKzBjTb5X\n+dXv/xzSfoxk7sbDqhPg4f7A1HDt7S3cUoDMMbPx1rxtGPDXZXhr3raAnx/8t2UYOXWtYyOKRgTU\nYokfP1OGX3fW5OKdnpuPJdsL8beF3vel575embUZq/e5ryhetLUAaz1WGctAq4a13baDcU9+fsYG\nfLfW+/dlNJYS9EVbC5D10nyv8q4Z5gl6MIkvSpUgQxN/DtxtDwbPlZ3B8JfvN2HU1+u9foCnwtiX\nUTzy5RpMWrEbr/+4BWUVDhH/Yf0htBs7F7sKgw8V4CluM9cdrBFhF/6hdI//tdQxibpq73E33+fl\neYX4aOnO6hY5AJRH6WKX/OPuInT6fAW6v74Ad0zKqX4gOi2fumo/8o/XxMRftLUAxS5BxqoE8O9f\n9+LWj3912+f9U3Jx80fuZZ6cPFuOvUVnqreFEF4P6QHvLsVD/8kN+tz0xvN+WbajEC9+vxE/bT7i\nVbeqSqDH6wswddUBPDVtvekeeJYS9EnL1dNBmelJ4HpoIURQQzDnKyq9bpqKyqrqz67aU4wf1h9C\nuUtL9PDJc0FF8lu9rxjtx87B0ZJSnDzniLpXUlqO79YeBGBOwCNf9G6T5radu7fYq874OdswecUe\nfL/OYf/cTYcBQHXI6cMlO5E5ZrZbr+nIyVIcOuk+pPDEV+vcRBlwDD2tP3DCfYce1+iuyavw9rzt\nqg8DALjhH7+olhtBnST3CfucPTXfZVlFFTq/XNMQcp6V6z34wBSHoJ4qLcf9U3Ld4qj88ZPflPq+\nj3/6fAUyx8xG5pjZ1ffp5BV70O21n9DPZVjr9R+3os0Lc1BZJXDybDn2HzuL3UVn8NOWgpDOV088\nzzN373F8sXI/Hvrv6uqyFXlFOHjiHM5XVLk9/MxuIFlK0AXMbw3VincP0/utyzDGwPeWocOL86q7\naL/uKsKeojNuVp88V44OL87D41PXurVcbpu4Eh1enIec3cfwx09+w+NT1+KtuTVd/t5vLkb31xfg\nxNmamydzzGxM/HkXeo1fVC3+nyzbjfJKgcveWoJur/4EALjTJURBoTJxd7asAndNNjd0gedz+L0F\nO7BsR6F6ZY9LTyDMWJOPhS5C8Hdl6MD1QTjovWXuu1FRpWm/70efCYv92viPxd6TW54PhSMmjkW3\nb+TuBvrMN+sBAB8v24ULX5oX8PNnlQe9c3m9894sKCn1GmoBgHnKg9WJa89v86ESfLxsl2r6ts9+\ncTTKXpm1Gd1e+wlXvLNE1Z4th6InnLLndQaA4ZNzcNVfl2GdZyNAhTUGDktZStBX7vZuwRnNS9dn\nuW2P/W4TCpQf8u5Cx49g3X7HRb7j0xz0f3cpypXhAgKqRfbHDYfdWi7OH81tE1dWl01ascfrxr7o\ntQVu2+PnbMORklIs2FKAXYUeKrHlAAATHUlEQVSnq1s6ZS6itvHgyerXT3y1DgCwbHshlucVwUxG\n9Ml02/511zGM+GyVal2nsDr1eMGWIxj19Xo8oHTVX565qfqcv1i5D0IICCFwysWL5aYPf0Hr5+d4\n7Xv0txtVj7k+/6RquRNfLXUz+Pyenl5lB0+cw4S527zmeZzPUddnmzMNYc33LLAirwiXjnd3QHA+\nEB/+Yk112cx1B3HX5JrrNnxyDibMdZ9/OFpSiuV5NQ9rZ4/Lk93KUNrQD0KbIzKSXsp3cq68Erd/\nutLr/WOnz6PLKzU9ot/3FGs67+MPSzhW7zx6Cku3+2i5GYxa4Ko+ExZj1/ih1duHTpxza/3ep3Rf\nT4URnfHWj3/FlteuCVjv0S/XqJYfPuk9oVNSWo7zFeanfEsLIZ7Ls9M3ICEuBnM3OcYxv/cIUzvl\ntxo3x3Gzt+LT5btRUOLuRlh0OvTkE+8t2KFanj3O/cE604dAGUW9ZO8QA3199Dpu/uhXbDx40itC\n6Yvfb8QDlzkibZaUVmC4Sg9uzf7juLiVe8xwZyPBHz3HL3KLPKk2NHHbJ78hZ0+xqXNiANC0XpLP\n90pKy/32xG78cAX+L7uF2/m9OXcbvs49gEVP99PSTFUsIehXvRc9nhhqo/XO8UAnntHuArX0/HG2\nrLJ6DDMcer/p/aPu+spPPusbGUwq1KkPX8KxW2WC1FPMtcbz4aBmW7QuEnX22Dx98L9YuR9frNyv\n9pFqpq8+WD3eHiqeD2FPnOP+GyL4vWiBv/jygVw89x47q+qquKvwjEpt7bGEoEcTLdOSVcu7veZb\nJCNl1R7jhpqevOoCw46llZ+fr96J2cgYHXDqKv+CLzuRNC4rq4Tuq10tNYYeDVySae0UVdFEpo+H\nY6hsO+KdTDgasHoG+WiiRYNaZpsQMWpJr7WGBZ0xjbSURLNN0JVLMuubbYI0qKwntBxG9NhY0BmG\niXo8/eytiBFu1yzoDKMTbQ0MESw7Azo2MtuEiOEWOsNYmFZptQNXYoLiUo9VxXrTMEX7FIks6Iyh\nZDXlrOpMdHLlBemGHq95qvaTsDzkEgQ9WqaabYI03NKjudkmMEx0oMMiAm6hB4HakmcmPBrX9b1C\njmGYyDBiWYLlBT0aAnbJQufm5i65ZphooWUDbdZIGI3lBd1iSWIYm9CztfEL0BL8LFlnQmPCLV00\n36dapE+tsfwdYMSXxOjH6zd2MtsEXWiQrL2XRCDG6yBCdqV2ovZ+71E/5EJE1xDRdiLaSURjtDIq\nFFjOrU1CnOXbFKq0amh8l/3m7jypHc1E9aQoEcUC+BDAEABZAG4noiz/n9IeMxronklkZaFFfePj\nZRgUJtpwnhncwfBj6h34iYmMaB9y6QlgpxBitxCiDMBXAG7UxqzgMWNStHW6nAtG/IUN1Qs9FnBE\nA2px8xl7E+1DLs0BHHDZzlfKDKVukndgf8Y68JAZYxeiesgF6sGsvUwmooeIKJeIcgsLtc86lOSR\n45OxFk3Y950JQN92xi7714toH3LJB9DCZTsDgFdKEiHERCFEthAiOz3d2OW7TPTTrQWv9GX845kq\nz6qopQnUmkgE/XcA7YmoNRElABgGYJY2ZkU3SXHcK2AYo3i0XztTjqt1/Ji02vrH/w9b0IUQFQD+\nDGA+gK0AvhZCbNbKsGimfgjJjRmGiQyzenF/uDhD0/3tL9Y/r2hEfRkhxBwAczSyhWEYJmrQ3gtU\nf7dS9q1iGIl4ytAk33Kj9RymEcsELC3oXTM4mBTDuHL5BQ3NNoHxAekQktcTSws6IwdpPCehGTEG\niIZd0NrJkFvoAeBbVw5kjediBrz6P3ohHkNnjOKePpmmHZvHfbXDCNGwC03rabvozYjOk6UFfVBW\nY7NNkIb7+rY27dg3XNTMtGPLhhlRHmXlkkxtY9obMRxmaUHv044ngLSCh17lQLbYRhkmRADVizYG\nBPWztKDHmqhCl5qQkYZh7IaZPUet4RZ6ABLjzTM/TdKwr0zkpBoQs8MuyPRd8hi6H8bf3AUdm9Q1\n7fhG+JTaBdlih7dKkzNevhl0aFLHbBM0wwjFsOQv6faeLXHHpS1NtaFvW+uO37dK8544a5iif+Ag\nX8TGEGonyBPwrHkqhwTWCplSBhuRQMaSgv7itReabQJu79kicCUDuDWMAEKeY3lN6iahlsmCOrx3\nq4g+v/SZfkHXTdJ5qK5Lcw4JrBVNNHYdNIuOBvU0LCHoX9x/KXpmNsCSZ/ph+7hrdMnIHSq+hlwS\ng1wk89J1Wfh97FUR2fDZPdl47cZOIX/O0/IWDcz3JHhmcIewfeFfvaETMhsGP8wRHxP5bf/x8B6q\n5fOfvAJ/uqJNxPs3gxYNamHLa1eH/fn4WG0HFZY928/UniMAzBl5udt2uJO0Rg3RWkLQL2vfEF8/\n3ButG9ZGYhTGIndNzrvxlZofRE8/fqz3XdYa6XVqbtbrujatfj3lvp7Vr//v4gzMHnkZ4lSWADap\nWwvJCXHYO+Ha6jLnj6qWv0xOHrvq1cb8jDDxsTF45Qb3h9ONKv7pag/MxkrWo5ED27uVPz+kY/Xr\nHeOG4NmrO+Cydg3x/NDIengJcTHo16GR6nsdmtRBjMnLNdUe0E94fDeejBzYHsufG4DkhDhMuKUL\n0monYHivliEJWN4bQz22h+APPQL3IC9s6j0XlpocHxWtc8/vctTgC6ozKI1QepUDOzZyGzVY+5dB\nXvsZ0rmJjlbWYAlBj3aeHuxY6Tg4q7HbMvavH+7tJrb+ePvWrgCA9o1S0MLF93bczZ3RqVk9jFUZ\nZmqeWlNvzsjLkfPCwOog+v405YZuzVA3ydHLmfpgLzwZRSs1b+lek5Y2OcG9J5aSGIft44bgm4d7\nV5el10nEwAsd4jq8V828yp/7t8NDLi3lhLgYPNa/Hb544FJ0ae4e1M3ZeFr+XH98dKd7y/uBy7wF\nbce4IUiKj8XCUVe4ld8d4bCRVjhXi76j3FOA+8PO82G/cNSVeNLl/WE9W2L1XwZh3E1d0Kdt4If9\n13/qjRWj+3uVx8fG4JLM+gE/7/yOXYfC1r00OCoab3WS4tG4bk3DK4a8E9wQEVKTHV5vt3Rv7pUv\n4bquTfHn/sYk6TB/7EIiWrssHFCbeIyLIVRUCQzt4v20Tk6Iw4rR/dGoThJOlZZXlztv6rt7ZyIp\nPhZfrNyHzYdKALintMpq5t7KcXbxHunXFrPWHcLBE+eq38tu1SCqRNyVt2/tihlrDwJwd/Oacl9P\ntGuUAsCxgu/qTo0xf3MBXruhU7WXTKM6SVj1wkAkxMVU/8DU6OISpTOGaoIwNUuthRYNktG9ZSrW\n7j+B5qm10MpjKKe1y3a7Ru7jomoPXTNxXeno2otcMOoKZNRPxvoDJ/DV7wfQNr22zyEB58MyITYG\nZZVVABy9pPMVVfj2kT6oVyvO7Xv48oFLceeknOrt2y5pge/WHkTOnmI8P6QjerZugIKSUhSUnMfL\nsxz5cJyNoFsvzsDwXq2w8+hpjb4BbVjyTD9kvTQfgGP+acIfuuKzX/ZgUFZjTPltHzo0SanJF6p8\njbkvXoXscQsBAC9f38mwXhsLegR8+0hvVFQKHDtTBgBoo/zYlz7TTzWr0c7xQ73KXMmo73gIpKmM\nG8bGEG7v2RK392yJzDGzfe7j+m5N8enyPfj7sIvwzvztePKq9hh9TUfM23QED3+xGkB0B3CKi43B\nrD/3RUWVQFrtBKzeexz/e/BSr+/E2Qr1dIJo5JF0+uXrs1ByrsLn8ZaPHoCnv16HlbuLq8sm3Z2N\n79YexO09W+LM+Qq8v2BH9TWe9qdeqvt5rH/bqGhRAjXX1/O7efbqDnh/4Y7q+6xbi9SA2YCICMue\n7Ye0lER0fnm+23tZTet6Tab39Vi9TUTVPcm0lER0b1nTYncK+tAuTbG36Azuvaw1UhLjTHVHViM5\nIQ5r/jIImw+dRFJ8LJLiYzH6Gsdw3reP9Ea3jNTqRojT4aBhSiLqJ8fj+Nlyt4ep3rCgR8DFrRwt\nICEEvn2kN3ooN6vnBN3YoRf6nLQbOaAdPli806t89sjLcOhEacg2jRlyIR4f2B51k+Ix8EL1WDdm\ne7QEomtGjcjMf+oK1Tr1azt6J4Emoe8NMAbcPLUWJo24BAeKz1b/8NJSEvHA5Y7hmtqJcVj9l0E4\nUHwWJ8+Vo1Ed9wfGLT2aY8aag1ExUe/E2dqu8vD5e6x/OzwWRtffl199sPN8DZTGja9kz7ExhMcD\njPGbTYPaCbi8vXeOUacGwL2BDgD46M6L8cnPu1CvlnGLo6LnLrQwRFRzYVV40I/Xw6jBHTBqcAev\n8k7N6qFTs9ATeMTGUMB4Hq6tJKsy9tostGtUBwM6qk9OhkJKYpzqxJwrLRokQ81RdezQC1Fyrhx3\n9oyO8XOgRlSEAAZ0bIRbejT3W19rFj19JfKP1wzxPXN1B7RJT8HVneQNptdZmZdxvR97t01D7yDm\nILSEBd0m9G2Xho5N6uD9YReZbYompCTG4X6VCUujSUtJxKQRl5hthhvO5fKxMYTP7tHOtg6N62B7\nwSk81r8d3luww+cK37bpKWibnlK9nRQfq7oQMCE2Bg9eYf411IKsZnWx7fVrkOTPu8wAWNAtSI+W\nqRjapWngii7USYrHvCfVhy8Yufh4+MWYs/Gw2wSuFkz7Uy8cKD6HLhn1vFxEw2HHG0M0sCp6MFvM\nARZ0SzLj0b5mm8BEMY3qJuEeHaIUpiYn+PUeYsyHBZ2xJV/cfymOnTlvthkMoyks6Iwtuay9dYOr\nMYwveKUowzCMJLCgMwzDSAILOsMwjCSwoDMMw0gCCzrDMIwksKAzDMNIAgs6wzCMJLCgMwzDSAIJ\nA9NqE1EhgH1hfrwhgCINzbECfM72gM9ZfiI931ZCCO/4vR4YKuiRQES5Qohss+0wEj5ne8DnLD9G\nnS8PuTAMw0gCCzrDMIwkWEnQJ5ptgAnwOdsDPmf5MeR8LTOGzjAMw/jHSi10hmEYxg+WEHQiuoaI\nthPRTiIaY7Y9oUBELYhoCRFtJaLNRPSEUt6AiBYQUZ7yv75STkT0gXKuG4ioh8u+Rij184hohEv5\nxUS0UfnMB0TB5mPXFyKKJaK1RPSjst2aiHIU+6cRUYJSnqhs71Tez3TZx/NK+XYiutqlPOruCSJK\nJaLpRLRNud69Zb/ORPSUcl9vIqKpRJQk23Umos+I6CgRbXIp0/26+jqGX4QQUf0HIBbALgBtACQA\nWA8gy2y7QrC/KYAeyus6AHYAyALwNoAxSvkYAG8pr4cCmAtH8vZeAHKU8gYAdiv/6yuv6yvvrQLQ\nW/nMXABDzD5vxa5RAP4H4Edl+2sAw5TXHwN4RHn9KICPldfDAExTXmcp1zsRQGvlPoiN1nsCwBQA\nDyivEwCkynydATQHsAdALZfre49s1xnAFQB6ANjkUqb7dfV1DL+2mv0jCOLL7A1gvsv28wCeN9uu\nCM5nJoBBALYDaKqUNQWwXXn9CYDbXepvV96/HcAnLuWfKGVNAWxzKXerZ+J5ZgBYBGAAgB+Vm7UI\nQJzndQUwH0Bv5XWcUo88r7WzXjTeEwDqKuJGHuXSXmc4BP2AIlJxynW+WsbrDCAT7oKu+3X1dQx/\nf1YYcnHeNE7ylTLLoXQxuwPIAdBYCHEYAJT/jZRqvs7XX3m+SrnZvA/gOQBVynYagBNCiApl29XO\n6nNT3j+p1A/1uzCTNgAKAXyuDDNNIqLakPg6CyEOAngXwH4Ah+G4bqsh93V2YsR19XUMn1hB0NXG\nCS3nmkNEKQC+BfCkEKLEX1WVMhFGuWkQ0XUAjgohVrsWq1QVAd6zzDnD0eLsAeBfQojuAM7A0U32\nheXPWRnTvRGOYZJmAGoDGKJSVabrHAhTz9EKgp4PoIXLdgaAQybZEhZEFA+HmH8phJihFBcQUVPl\n/aYAjirlvs7XX3mGSrmZ9AVwAxHtBfAVHMMu7wNIJSJnYnJXO6vPTXm/HoBihP5dmEk+gHwhRI6y\nPR0OgZf5Ol8FYI8QolAIUQ5gBoA+kPs6OzHiuvo6hk+sIOi/A2ivzJwnwDGZMstkm4JGmbGeDGCr\nEOI9l7dmAXDOdI+AY2zdWX63MlveC8BJpbs1H8BgIqqvtIwGwzG+eBjAKSLqpRzrbpd9mYIQ4nkh\nRIYQIhOO67VYCHEngCUAblWqeZ6z87u4VakvlPJhindEawDt4ZhAirp7QghxBMABIuqgFA0EsAUS\nX2c4hlp6EVGyYpPznKW9zi4YcV19HcM3Zk6qhDAhMRQO75BdAMaabU+Itl8GRxdqA4B1yt9QOMYO\nFwHIU/43UOoTgA+Vc90IINtlX/cB2Kn83etSng1gk/KZf8JjYs7k8++HGi+XNnD8UHcC+AZAolKe\npGzvVN5v4/L5scp5bYeLV0c03hMALgKQq1zr7+HwZpD6OgN4FcA2xa7/wuGpItV1BjAVjjmCcjha\n1PcbcV19HcPfH68UZRiGkQQrDLkwDMMwQcCCzjAMIwks6AzDMJLAgs4wDCMJLOgMwzCSwILOMAwj\nCSzoDMMwksCCzjAMIwn/D1/lqqhe580EAAAAAElFTkSuQmCC\n",
      "text/plain": [
       "<matplotlib.figure.Figure at 0x7f582f0e2dd8>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%matplotlib inline\n",
    "from matplotlib import pyplot as plt\n",
    "plt.plot(waveform)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 评估\n",
    "\n",
    "下面的练习将用到您目前所学的全部知识。不同于之前的练习，本次练习不提供任何解决方案代码。就像本节一样，本课程中的其它2个部分的notebook里也有需要评估的问题。成功完成全部课程的评估问题后，您将获得本课程的**“能力证书”**。\n",
    "\n",
    "**请仔细阅读说明后再开始工作，确保以最大几率成功完成本次评估。**"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 加速神经网络计算\n",
    "\n",
    "下例简单演示了在神经网络中创建隐藏层所需执行的部分工作。该示例对一百万个灰度值（此处随机创建）执行归一化，为其赋予权重，然后应用激活函数对其作出处理。\n",
    "\n",
    "您要完成的任务是：运用所学技巧将此工作移至 GPU，确保正确执行计算，并提升函数调用性能。\n",
    "\n",
    "开始任务之前，请注意以下事项：\n",
    "\n",
    "* 仅向设备发送一次值，并将其保留在设备上以用于多次运算。相较于每次函数调用都在主机和设备之间来回传送数据，此举将能显著提升性能。本课程中，您已经学习如何将设备数组用作输出，以及如何消减 GPU 和主机间的数据往返，进而在 GPU 上重新使用数组。\n",
    "* `np` 数学函数无法在设备上使用，您需改用与其作用相同的 `math` 函数。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 导入库并进行初始化"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "运行此单元格以导入所需的库并初始化值，然后再开始下面的工作。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 94,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "# You should not modify this cell, it contains imports and initial values needed to do work on either\n",
    "# the CPU or the GPU.\n",
    "\n",
    "import numpy as np\n",
    "from numba import cuda, vectorize\n",
    "\n",
    "# Our hidden layer will contain 1M neurons.\n",
    "# When you assess your work below, this value will be automatically set to 100M.\n",
    "n = 1000000\n",
    "\n",
    "greyscales = np.floor(np.random.uniform(0, 255, n).astype(np.float32))\n",
    "weights = np.random.normal(.5, .1, n).astype(np.float32)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 用GPU加速"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "在评估下面的工作之前，您需要对本节的3个单元中的每个单元进行修改。请按照注释中的说明进行操作。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 95,
   "metadata": {},
   "outputs": [],
   "source": [
    "# As you will recall, `numpy.exp` works on the CPU, but, cannot be used in GPU implmentations.\n",
    "# This import will work for the CPU-only boilerplate code provided below, but\n",
    "# you will need to modify this import before your GPU implementation will work.\n",
    "# from numpy import exp\n",
    "from math import exp\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 96,
   "metadata": {},
   "outputs": [],
   "source": [
    "# Modify these 3 function calls to run on the GPU.\n",
    "@vectorize(['float32(float32)'], target='cuda')\n",
    "def normalize(grayscales):\n",
    "    return grayscales / 255\n",
    "\n",
    "@vectorize(['float32(float32, float32)'], target='cuda')\n",
    "def weigh(values, weights):\n",
    "    return values * weights\n",
    "\n",
    "@vectorize(['float32(float32)'], target='cuda')\n",
    "def activate(values):\n",
    "    return ( exp(values) - exp(-values) ) / ( exp(values) + exp(-values) )"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 109,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "# Modify the body of this function to optimize data transfers and therefore speed up performance.\n",
    "# As a constraint, even after you move work to the GPU, make this function return a host array.\n",
    "\n",
    "def create_hidden_layer(n, greyscales, weights, exp, normalize, weigh, activate):\n",
    "    \n",
    "    d_greyscales = cuda.to_device(greyscales)\n",
    "    d_weights = cuda.to_device(weights)\n",
    "    d_normalized = cuda.device_array(shape=(n,), dtype=np.float32)\n",
    "    d_activated = cuda.device_array(shape=(n,), dtype=np.float32)\n",
    "    d_new_weights = cuda.device_array(shape=(n,), dtype=np.float32)\n",
    "    \n",
    "    normalize(d_greyscales,out=d_normalized)\n",
    "    weigh(d_normalized, d_weights,out=d_new_weights)\n",
    "    activate(d_new_weights,out=d_activated)\n",
    "    \n",
    "    # The assessment mechanism will expect `activated` to be a host array, so,\n",
    "    # even after you refactor this code to run on the GPU, make sure to explicitly copy\n",
    "    # `activated` back to the host.\n",
    "    return d_activated.copy_to_host()"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 检查您的工作"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "在运行以下评估之前，可以随时在本节中检查您的工作并根据需要进行调试。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 110,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "# You probably don't need to edit this cell, unless you change the name of any of the values being passed as\n",
    "# arguments to `create_hidden_layer` below.\n",
    "arguments = {\"n\":n,\n",
    "            \"greyscales\": greyscales,\n",
    "            \"weights\": weights,\n",
    "            \"exp\": exp,\n",
    "            \"normalize\": normalize,\n",
    "            \"weigh\": weigh,\n",
    "            \"activate\": activate}"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 111,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "[0.17218623 0.05202709 0.12721425 ... 0.200005   0.1778189  0.14301781]\n"
     ]
    }
   ],
   "source": [
    "# Use this cell (and feel free to create others) to self-assess your function\n",
    "a = create_hidden_layer(**arguments)\n",
    "print(a)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 运行评估"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "运行以下2个单元格以评估您的工作。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 112,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "from assessment import assess"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 113,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Setting n to 100 million.\n",
      "\n",
      "Your function returns a host np.ndarray: True\n",
      "\n",
      "Your function took 0.47s to run.\n",
      "Your function runs fast enough (less than 1 second): True\n",
      "\n",
      "Your function returns the correct results: True\n",
      "Congratulations, you passed! See the instructions below for how to get credit for your work to count toward a certificate in the course.\n"
     ]
    }
   ],
   "source": [
    "assess(create_hidden_layer, arguments)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 为您的工作获得评估分数"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "成功通过上述评估后，请重新访问启动此交互式环境的网页，然后单击**“ASSESS TASK”** 按钮，如下面的截屏所示。这样将使您获得这部分课程的得分，它是您获得整个课程的**“能力证书”**的一部分。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![Run the assessment](images/run_the_assessment.png)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 总结\n",
    "\n",
    "您已完成本课程学习，现已能够：\n",
    "\n",
    "- 使用 Numba 为 CPU 编译 Python 函数\n",
    "- 了解 Numba 如何编译函数\n",
    "- 使用 GPU 加速 NumPy ufunc\n",
    "- 使用 GPU 加速手写向量化函数\n",
    "- 优化 CPU 主机和 GPU 设备之间的内存传输"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 下载内容\n",
    "\n",
    "如要下载此笔记本的内容，请执行以下单元，然后点击下方的下载链接。注意：由于笔记本中的部分文件路径链接是专为我们的平台量身设计，若您在本地 Jupyter 服务器上运行此笔记本，这些链接可能会遭到损坏。不过，您仍可通过 Jupyter 文件导航器导航至这些文件。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 114,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "./\r\n",
      "./.ipynb_checkpoints/\r\n",
      "./.ipynb_checkpoints/Introduction to CUDA Python with Numba-checkpoint.ipynb\r\n",
      "./Introduction to CUDA Python with Numba.ipynb\r\n",
      "./solutions/\r\n",
      "./solutions/make_pulses_solution.py\r\n",
      "./solutions/zero_suppress_solution.py\r\n",
      "./solutions/monte_carlo_pi_solution.py\r\n",
      "./images/\r\n",
      "./images/DLI Header.png\r\n",
      "./images/run_the_assessment.png\r\n",
      "./images/numba_flowchart.png\r\n",
      "tar: .: file changed as we read it\r\n"
     ]
    }
   ],
   "source": [
    "!tar -zcvf section1.tar.gz ."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "[下载本节文件。](files/section1.tar.gz)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 附录：广义 ufunc\n",
    "\n",
    "ufunc 可针对数组输入广播标量函数，但若要针对高维数组广播低维数组函数，又该如何操作呢？此时我们便会用到名为*广义 ufunc* (\"gufunc\") 的函数，这为应用 ufunc 开拓了一片新天地。\n",
    "\n",
    "广义 ufunc 相对来说更加复杂，因为此类函数在处理多个输入时，需要使用*签名*（请勿与 Numba 类型签名混淆）来显示索引排序。全面讲解 \"gufunc\" 签名已超出本教程大纲，但您可参阅以下文件以作深入了解：\n",
    "\n",
    "* gufunc 之 NumPy 讲解文档： https://numpy.org/neps/nep-0005-generalized-ufuncs.html\n",
    "* gufunc 之 Numba 讲解文档：http://numba.pydata.org/numba-doc/latest/user/vectorize.html#the-guvectorize-decorator\n",
    "* CUDA gufunc 之 Numba 讲解文档：http://numba.pydata.org/numba-doc/latest/cuda/ufunc.html#generalized-cuda-ufuncs\n",
    "\n",
    "下面我们就来编写自己的归一化函数。编写时需要提供数组输入，并需沿最后一个维度计算 L2 范数。广义 ufunc 会将输出数组用作最后一个参数，而非返回值。若输出为标量，则我们仍可获得一个数组，但该数组会比输入数组少一个维度。例如，在计算数组的行总和时，输入二维数组将会返回一维数组，而输入三维数组则会返回二维数组。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "from numba import guvectorize\n",
    "import math\n",
    "\n",
    "@guvectorize(['(float32[:], float32[:])'], # have to include the output array in the type signature\n",
    "             '(i)->()',                 # map a 1D array to a scalar output\n",
    "             target='cuda')\n",
    "def l2_norm(vec, out):\n",
    "    acc = 0.0\n",
    "    for value in vec:\n",
    "        acc += value**2\n",
    "    out[0] = math.sqrt(acc)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "如要测试这一点，我们需在单位圆中绘制一些点："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "angles = np.random.uniform(-np.pi, np.pi, 10)\n",
    "coords = np.stack([np.cos(angles), np.sin(angles)], axis=1)\n",
    "print(coords)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "一如预期，在不考虑舍入误差时，L2 范数为 1.0："
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "l2_norm(coords)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "<a href=\"https://www.nvidia.com/dli\"> <img src=\"images/DLI Header.png\" alt=\"标题\" style=\"width: 400px;\"/> </a>"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.6.10"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 2
}
